Skip to content

Commit 072a8f7

Browse files
mylermyler
authored
Merge pull request intel#946 from bb-sycl/xmain
Auto pulldown and update tc files for xmain branch on 20220530
2 parents 170cb56 + 054e93b commit 072a8f7

17 files changed

+87
-133
lines changed

‎SYCL/Basic/partition_supported.cpp

Lines changed: 30 additions & 26 deletions
Original file line numberDiff line numberDiff line change
@@ -6,15 +6,13 @@
66
// Nvidia should not allow sub_devices but does not throw corresponding error.
77
// XFAIL: hip_nvidia
88
/* Check that:
9-
1) [info::device::partition_properties]: returns the partition properties
10-
supported by this SYCL device; a vector of info::partition_property. If this
11-
SYCL device cannot be partitioned into at least two sub devices then the
12-
returned vector **must be empty**.
9+
1) if partition_equally is supported, then we check that the correct
10+
invalid errc is returned if more than max_compute_units are requested
1311
14-
2) [create_sub_devices()]: If the SYCL device
15-
does not support info::partition_property::partition_by_affinity_domain or the
16-
SYCL device does not support the info::partition_affinity_domain provided, an
17-
exception with the **feature_not_supported error code must be thrown**.
12+
2) If the SYCL device does not support
13+
info::partition_property::partition_by_affinity_domain or the SYCL device does
14+
not support the info::partition_affinity_domain provided, an exception with the
15+
**feature_not_supported error code must be thrown**.
1816
*/
1917

2018
#include <CL/sycl.hpp>
@@ -57,24 +55,10 @@ int main() {
5755

5856
auto dev = cl::sycl::device(cl::sycl::default_selector());
5957

60-
cl::sycl::info::partition_property partitionProperty =
61-
cl::sycl::info::partition_property::partition_by_affinity_domain;
62-
cl::sycl::info::partition_affinity_domain affinityDomain =
63-
cl::sycl::info::partition_affinity_domain::next_partitionable;
64-
65-
if (supports_partition_property(dev, partitionProperty)) {
66-
if (supports_affinity_domain(dev, partitionProperty, affinityDomain)) {
67-
auto subDevices = dev.create_sub_devices<
68-
cl::sycl::info::partition_property::partition_by_affinity_domain>(
69-
affinityDomain);
70-
71-
if (subDevices.size() < 2) {
72-
std::cerr << "device::create_sub_device(info::partition_affinity_"
73-
"domain) should have returned at least 2 devices"
74-
<< std::endl;
75-
return -1;
76-
}
77-
}
58+
// 1 - check exceed max_compute_units
59+
cl::sycl::info::partition_property partitionEqually =
60+
cl::sycl::info::partition_property::partition_equally;
61+
if (supports_partition_property(dev, partitionEqually)) {
7862
auto maxUnits = dev.get_info<sycl::info::device::max_compute_units>();
7963
try {
8064
std::vector<sycl::device> v = dev.create_sub_devices<
@@ -93,6 +77,26 @@ int main() {
9377
return -1;
9478
}
9579
}
80+
}
81+
82+
// 2 - check affinity
83+
cl::sycl::info::partition_property partitionProperty =
84+
cl::sycl::info::partition_property::partition_by_affinity_domain;
85+
cl::sycl::info::partition_affinity_domain affinityDomain =
86+
cl::sycl::info::partition_affinity_domain::next_partitionable;
87+
if (supports_partition_property(dev, partitionProperty)) {
88+
if (supports_affinity_domain(dev, partitionProperty, affinityDomain)) {
89+
auto subDevices = dev.create_sub_devices<
90+
cl::sycl::info::partition_property::partition_by_affinity_domain>(
91+
affinityDomain);
92+
93+
if (subDevices.size() < 2) {
94+
std::cerr << "device::create_sub_device(info::partition_affinity_"
95+
"domain) should have returned at least 2 devices"
96+
<< std::endl;
97+
return -1;
98+
}
99+
}
96100
} else {
97101
try {
98102
auto subDevices = dev.create_sub_devices<

‎SYCL/GroupAlgorithm/SYCL2020/exclusive_scan.cpp

Lines changed: 3 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -3,14 +3,9 @@
33
// RUN: %GPU_RUN_PLACEHOLDER %t.out
44
// RUN: %ACC_RUN_PLACEHOLDER %t.out
55

6-
// TODO: enable compile+runtime checks for operations defined in SPIR-V 1.3.
7-
// That requires either adding a switch to clang (-spirv-max-version=1.3) or
8-
// raising the spirv version from 1.1. to 1.3 for spirv translator
9-
// unconditionally. Using operators specific for spirv 1.3 and higher with
10-
// -spirv-max-version=1.1 being set by default causes assert/check fails
11-
// in spirv translator.
12-
// RUNx: %clangxx -fsycl -fsycl-targets=%sycl_triple -DSPIRV_1_3 %s -I . -o \
13-
%t13.out
6+
// disabling hip because some of the binary_ops tested are not supported
7+
// getting undefined symbols for a handful of __spirv__ * functions.
8+
// XFAIL: hip
149

1510
#include "support.h"
1611
#include <CL/sycl.hpp>
@@ -149,16 +144,13 @@ int main() {
149144
std::numeric_limits<int>::max());
150145
test<class KernelNameMaximumI>(q, input, output, sycl::maximum<int>(),
151146
std::numeric_limits<int>::lowest());
152-
153-
#ifdef SPIRV_1_3
154147
test<class KernelName_VzAPutpBRRJrQPB>(q, input, output,
155148
sycl::multiplies<int>(), 1);
156149
test<class KernelName_UXdGbr>(q, input, output, sycl::bit_or<int>(), 0);
157150
test<class KernelName_saYaodNyJknrPW>(q, input, output, sycl::bit_xor<int>(),
158151
0);
159152
test<class KernelName_GPcuAlvAOjrDyP>(q, input, output, sycl::bit_and<int>(),
160153
~0);
161-
#endif // SPIRV_1_3
162154

163155
std::cout << "Test passed." << std::endl;
164156
}

‎SYCL/GroupAlgorithm/SYCL2020/inclusive_scan.cpp

Lines changed: 3 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -3,14 +3,9 @@
33
// RUN: %GPU_RUN_PLACEHOLDER %t.out
44
// RUN: %ACC_RUN_PLACEHOLDER %t.out
55

6-
// TODO: enable compile+runtime checks for operations defined in SPIR-V 1.3.
7-
// That requires either adding a switch to clang (-spirv-max-version=1.3) or
8-
// raising the spirv version from 1.1. to 1.3 for spirv translator
9-
// unconditionally. Using operators specific for spirv 1.3 and higher with
10-
// -spirv-max-version=1.1 being set by default causes assert/check fails
11-
// in spirv translator.
12-
// RUNx: %clangxx -fsycl -fsycl-targets=%sycl_triple -DSPIRV_1_3 %s -I . -o \
13-
%t13.out
6+
// disabling hip because some of the binary_ops tested are not supported
7+
// getting undefined symbols for a handful of __spirv__ * functions.
8+
// XFAIL: hip
149

1510
#include "support.h"
1611
#include <CL/sycl.hpp>
@@ -149,8 +144,6 @@ int main() {
149144
std::numeric_limits<int>::max());
150145
test<class KernelNameMaximumI>(q, input, output, sycl::maximum<int>(),
151146
std::numeric_limits<int>::lowest());
152-
153-
#ifdef SPIRV_1_3
154147
test<class KernelName_zMyjxUrBgeUGoxmDwhvJ>(q, input, output,
155148
sycl::multiplies<int>(), 1);
156149
test<class KernelName_SljjtroxNRaAXoVnT>(q, input, output,
@@ -159,7 +152,6 @@ int main() {
159152
sycl::bit_xor<int>(), 0);
160153
test<class KernelName_xGnAnMYHvqekCk>(q, input, output, sycl::bit_and<int>(),
161154
~0);
162-
#endif // SPIRV_1_3
163155

164156
std::cout << "Test passed." << std::endl;
165157
}

‎SYCL/GroupAlgorithm/SYCL2020/reduce.cpp

Lines changed: 3 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -3,14 +3,9 @@
33
// RUN: %GPU_RUN_PLACEHOLDER %t.out
44
// RUN: %ACC_RUN_PLACEHOLDER %t.out
55

6-
// TODO: enable compile+runtime checks for operations defined in SPIR-V 1.3.
7-
// That requires either adding a switch to clang (-spirv-max-version=1.3) or
8-
// raising the spirv version from 1.1. to 1.3 for spirv translator
9-
// unconditionally. Using operators specific for spirv 1.3 and higher with
10-
// -spirv-max-version=1.1 being set by default causes assert/check fails
11-
// in spirv translator.
12-
// RUNx: %clangxx -fsycl -fsycl-targets=%sycl_triple -DSPIRV_1_3 %s -I . -o \
13-
%t13.out
6+
// disabling hip because some of the binary_ops tested are not supported
7+
// getting undefined symbols for a handful of __spirv__ * functions.
8+
// XFAIL: hip
149

1510
#include "support.h"
1611
#include <CL/sycl.hpp>
@@ -86,15 +81,13 @@ int main() {
8681
test<class KernelNameMaximumI>(q, input, output, sycl::maximum<int>(),
8782
std::numeric_limits<int>::lowest());
8883

89-
#ifdef SPIRV_1_3
9084
test<class KernelName_WonwuUVPUPOTKRKIBtT>(q, input, output,
9185
sycl::multiplies<int>(), 1);
9286
test<class KernelName_qYBaJDZTMGkdIwD>(q, input, output, sycl::bit_or<int>(),
9387
0);
9488
test<class KernelName_eLSFt>(q, input, output, sycl::bit_xor<int>(), 0);
9589
test<class KernelName_uFhJnxSVhNAiFPTG>(q, input, output,
9690
sycl::bit_and<int>(), ~0);
97-
#endif // SPIRV_1_3
9891

9992
std::cout << "Test passed." << std::endl;
10093
}

‎SYCL/GroupAlgorithm/exclusive_scan.cpp

Lines changed: 3 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -4,14 +4,9 @@
44
// RUN: %GPU_RUN_PLACEHOLDER %t.out
55
// RUN: %ACC_RUN_PLACEHOLDER %t.out
66

7-
// TODO: enable compile+runtime checks for operations defined in SPIR-V 1.3.
8-
// That requires either adding a switch to clang (-spirv-max-version=1.3) or
9-
// raising the spirv version from 1.1. to 1.3 for spirv translator
10-
// unconditionally. Using operators specific for spirv 1.3 and higher with
11-
// -spirv-max-version=1.1 being set by default causes assert/check fails
12-
// in spirv translator.
13-
// RUNx: %clangxx -fsycl -fsycl-targets=%sycl_triple -DSPIRV_1_3 %s -I . -o \
14-
%t13.out
7+
// disabling hip because some of the binary_ops tested are not supported
8+
// getting undefined symbols for a handful of __spirv__ * functions.
9+
// XFAIL: hip
1510

1611
#include "support.h"
1712
#include <CL/sycl.hpp>
@@ -151,8 +146,6 @@ int main() {
151146
std::numeric_limits<int>::max());
152147
test<class KernelNameMaximumI>(q, input, output, ext::oneapi::maximum<int>(),
153148
std::numeric_limits<int>::lowest());
154-
155-
#ifdef SPIRV_1_3
156149
test<class KernelName_VzAPutpBRRJrQPB>(q, input, output,
157150
ext::oneapi::multiplies<int>(), 1);
158151
test<class KernelName_UXdGbr>(q, input, output, ext::oneapi::bit_or<int>(),
@@ -161,7 +154,6 @@ int main() {
161154
ext::oneapi::bit_xor<int>(), 0);
162155
test<class KernelName_GPcuAlvAOjrDyP>(q, input, output,
163156
ext::oneapi::bit_and<int>(), ~0);
164-
#endif // SPIRV_1_3
165157

166158
std::cout << "Test passed." << std::endl;
167159
}

‎SYCL/GroupAlgorithm/exclusive_scan_over_group.cpp

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -3,6 +3,9 @@
33
// RUN: %GPU_RUN_PLACEHOLDER %t.out
44
// UNSUPPORTED: ze_debug4,ze_debug-1
55

6+
// CPU and ACC not yet supported:
7+
// Unsupported SPIR-V module SPIRV module requires unsupported capability 6400
8+
69
#include <CL/sycl.hpp>
710
#include <algorithm>
811
#include <iostream>
@@ -36,10 +39,7 @@ int main(int argc, const char **argv) {
3639
int num_wg = 1;
3740
int group_size = 16;
3841

39-
cl::sycl::queue queue{
40-
cl::sycl::gpu_selector{},
41-
cl::sycl::property_list{cl::sycl::property::queue::enable_profiling(),
42-
cl::sycl::property::queue::in_order()}};
42+
cl::sycl::queue queue;
4343

4444
typedef int T;
4545
size_t nelems = num_wg * group_size;

‎SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp

Lines changed: 10 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -4,8 +4,6 @@
44
// RUN: %GPU_RUN_PLACEHOLDER %t.out
55
// RUN: %ACC_RUN_PLACEHOLDER %t.out
66

7-
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsycl-device-code-split=per_kernel -DSPIRV_1_3 %s -I . -o %t13.out
8-
97
#include "support.h"
108
#include <CL/sycl.hpp>
119
#include <algorithm>
@@ -45,7 +43,8 @@ void test(queue q, InputContainer input, OutputContainer output,
4543
typedef typename OutputContainer::value_type OutputT;
4644
typedef class exclusive_scan_kernel<SpecializationKernelName, 0> kernel_name0;
4745
constexpr size_t G = 64;
48-
constexpr size_t N = input.size();
46+
constexpr size_t N = input.size(); // 128 or 12
47+
constexpr size_t confirmRange = std::min(G, N);
4948
std::vector<OutputT> expected(N);
5049

5150
// checking
@@ -64,9 +63,10 @@ void test(queue q, InputContainer input, OutputContainer output,
6463
});
6564
});
6665
}
67-
emu::exclusive_scan(input.begin(), input.begin() + G, expected.begin(),
68-
identity, binary_op);
69-
assert(std::equal(output.begin(), output.begin() + G, expected.begin()));
66+
emu::exclusive_scan(input.begin(), input.begin() + confirmRange,
67+
expected.begin(), identity, binary_op);
68+
assert(std::equal(output.begin(), output.begin() + confirmRange,
69+
expected.begin()));
7070

7171
typedef class exclusive_scan_kernel<SpecializationKernelName, 1> kernel_name1;
7272
constexpr OutputT init = 42;
@@ -88,9 +88,10 @@ void test(queue q, InputContainer input, OutputContainer output,
8888
});
8989
});
9090
}
91-
emu::exclusive_scan(input.begin(), input.begin() + G, expected.begin(), init,
92-
binary_op);
93-
assert(std::equal(output.begin(), output.begin() + G, expected.begin()));
91+
emu::exclusive_scan(input.begin(), input.begin() + confirmRange,
92+
expected.begin(), init, binary_op);
93+
assert(std::equal(output.begin(), output.begin() + confirmRange,
94+
expected.begin()));
9495

9596
typedef class exclusive_scan_kernel<SpecializationKernelName, 2> kernel_name2;
9697

@@ -176,15 +177,12 @@ int main() {
176177
std::numeric_limits<int>::max());
177178
test<class KernelNameMaximumI>(q, input, output, sycl::maximum<int>(),
178179
std::numeric_limits<int>::lowest());
179-
180-
#ifdef SPIRV_1_3
181180
test<class KernelNameMultipliesI>(q, input_small, output_small,
182181
sycl::multiplies<int>(), 1);
183182
test<class KernelNameBitOrI>(q, input, output, sycl::bit_or<int>(), 0);
184183
test<class KernelNameBitXorI>(q, input, output, sycl::bit_xor<int>(), 0);
185184
test<class KernelNameBitAndI>(q, input_small, output_small,
186185
sycl::bit_and<int>(), ~0);
187-
#endif // SPIRV_1_3
188186

189187
// as part of SYCL_EXT_ONEAPI_COMPLEX_ALGORITHMS (
190188
// https://github.com/intel/llvm/pull/5108/ ) joint_exclusive_scan and

‎SYCL/GroupAlgorithm/inclusive_scan.cpp

Lines changed: 3 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -4,14 +4,9 @@
44
// RUN: %GPU_RUN_PLACEHOLDER %t.out
55
// RUN: %ACC_RUN_PLACEHOLDER %t.out
66

7-
// TODO: enable compile+runtime checks for operations defined in SPIR-V 1.3.
8-
// That requires either adding a switch to clang (-spirv-max-version=1.3) or
9-
// raising the spirv version from 1.1. to 1.3 for spirv translator
10-
// unconditionally. Using operators specific for spirv 1.3 and higher with
11-
// -spirv-max-version=1.1 being set by default causes assert/check fails
12-
// in spirv translator.
13-
// RUNx: %clangxx -fsycl -fsycl-targets=%sycl_triple -DSPIRV_1_3 %s -I . -o \
14-
%t13.out
7+
// disabling hip because some of the binary_ops tested are not supported
8+
// getting undefined symbols for a handful of __spirv__ * functions.
9+
// XFAIL: hip
1510

1611
#include "support.h"
1712
#include <CL/sycl.hpp>
@@ -151,8 +146,6 @@ int main() {
151146
std::numeric_limits<int>::max());
152147
test<class KernelNameMaximumI>(q, input, output, ext::oneapi::maximum<int>(),
153148
std::numeric_limits<int>::lowest());
154-
155-
#ifdef SPIRV_1_3
156149
test<class KernelName_zMyjxUrBgeUGoxmDwhvJ>(
157150
q, input, output, ext::oneapi::multiplies<int>(), 1);
158151
test<class KernelName_SljjtroxNRaAXoVnT>(q, input, output,
@@ -161,7 +154,6 @@ int main() {
161154
ext::oneapi::bit_xor<int>(), 0);
162155
test<class KernelName_xGnAnMYHvqekCk>(q, input, output,
163156
ext::oneapi::bit_and<int>(), ~0);
164-
#endif // SPIRV_1_3
165157

166158
std::cout << "Test passed." << std::endl;
167159
}

0 commit comments

Comments
 (0)