Skip to content
This repository was archived by the owner on Mar 28, 2023. It is now read-only.

Commit 46354b5

Browse files
cperkinsintelcperkinsintel
authored
[SYCL] Reenable GroupAlgorithm tests (#1021)
Some of the group algorithm operations had their full test coverage deferred when we were only supporting early versions of SPIR-V. We've long since moved on, yet the tests remained truncate. Here I am reenabling them to prevent any further regressions sneaking in. Note that there are a couple cases commented out because of L0 failures. I haven't investigated that yet, I'm merely opening a draft PR to make sure this first enablement is working as intended. Signed-off-by: Chris Perkins <[email protected]>
1 parent 139fdc9 commit 46354b5

10 files changed

+24
-103
lines changed

‎SYCL/GroupAlgorithm/SYCL2020/exclusive_scan.cpp

Lines changed: 0 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -3,15 +3,6 @@
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
14-
156
#include "support.h"
167
#include <CL/sycl.hpp>
178
#include <algorithm>
@@ -149,16 +140,13 @@ int main() {
149140
std::numeric_limits<int>::max());
150141
test<class KernelNameMaximumI>(q, input, output, sycl::maximum<int>(),
151142
std::numeric_limits<int>::lowest());
152-
153-
#ifdef SPIRV_1_3
154143
test<class KernelName_VzAPutpBRRJrQPB>(q, input, output,
155144
sycl::multiplies<int>(), 1);
156145
test<class KernelName_UXdGbr>(q, input, output, sycl::bit_or<int>(), 0);
157146
test<class KernelName_saYaodNyJknrPW>(q, input, output, sycl::bit_xor<int>(),
158147
0);
159148
test<class KernelName_GPcuAlvAOjrDyP>(q, input, output, sycl::bit_and<int>(),
160149
~0);
161-
#endif // SPIRV_1_3
162150

163151
std::cout << "Test passed." << std::endl;
164152
}

‎SYCL/GroupAlgorithm/SYCL2020/inclusive_scan.cpp

Lines changed: 0 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -3,15 +3,6 @@
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
14-
156
#include "support.h"
167
#include <CL/sycl.hpp>
178
#include <algorithm>
@@ -149,8 +140,6 @@ int main() {
149140
std::numeric_limits<int>::max());
150141
test<class KernelNameMaximumI>(q, input, output, sycl::maximum<int>(),
151142
std::numeric_limits<int>::lowest());
152-
153-
#ifdef SPIRV_1_3
154143
test<class KernelName_zMyjxUrBgeUGoxmDwhvJ>(q, input, output,
155144
sycl::multiplies<int>(), 1);
156145
test<class KernelName_SljjtroxNRaAXoVnT>(q, input, output,
@@ -159,7 +148,6 @@ int main() {
159148
sycl::bit_xor<int>(), 0);
160149
test<class KernelName_xGnAnMYHvqekCk>(q, input, output, sycl::bit_and<int>(),
161150
~0);
162-
#endif // SPIRV_1_3
163151

164152
std::cout << "Test passed." << std::endl;
165153
}

‎SYCL/GroupAlgorithm/SYCL2020/reduce.cpp

Lines changed: 0 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -3,15 +3,6 @@
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
14-
156
#include "support.h"
167
#include <CL/sycl.hpp>
178
#include <algorithm>
@@ -86,15 +77,13 @@ int main() {
8677
test<class KernelNameMaximumI>(q, input, output, sycl::maximum<int>(),
8778
std::numeric_limits<int>::lowest());
8879

89-
#ifdef SPIRV_1_3
9080
test<class KernelName_WonwuUVPUPOTKRKIBtT>(q, input, output,
9181
sycl::multiplies<int>(), 1);
9282
test<class KernelName_qYBaJDZTMGkdIwD>(q, input, output, sycl::bit_or<int>(),
9383
0);
9484
test<class KernelName_eLSFt>(q, input, output, sycl::bit_xor<int>(), 0);
9585
test<class KernelName_uFhJnxSVhNAiFPTG>(q, input, output,
9686
sycl::bit_and<int>(), ~0);
97-
#endif // SPIRV_1_3
9887

9988
std::cout << "Test passed." << std::endl;
10089
}

‎SYCL/GroupAlgorithm/exclusive_scan.cpp

Lines changed: 0 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -4,15 +4,6 @@
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
15-
167
#include "support.h"
178
#include <CL/sycl.hpp>
189
#include <algorithm>
@@ -151,8 +142,6 @@ int main() {
151142
std::numeric_limits<int>::max());
152143
test<class KernelNameMaximumI>(q, input, output, ext::oneapi::maximum<int>(),
153144
std::numeric_limits<int>::lowest());
154-
155-
#ifdef SPIRV_1_3
156145
test<class KernelName_VzAPutpBRRJrQPB>(q, input, output,
157146
ext::oneapi::multiplies<int>(), 1);
158147
test<class KernelName_UXdGbr>(q, input, output, ext::oneapi::bit_or<int>(),
@@ -161,7 +150,6 @@ int main() {
161150
ext::oneapi::bit_xor<int>(), 0);
162151
test<class KernelName_GPcuAlvAOjrDyP>(q, input, output,
163152
ext::oneapi::bit_and<int>(), ~0);
164-
#endif // SPIRV_1_3
165153

166154
std::cout << "Test passed." << std::endl;
167155
}

‎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: 0 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -4,15 +4,6 @@
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
15-
167
#include "support.h"
178
#include <CL/sycl.hpp>
189
#include <algorithm>
@@ -151,8 +142,6 @@ int main() {
151142
std::numeric_limits<int>::max());
152143
test<class KernelNameMaximumI>(q, input, output, ext::oneapi::maximum<int>(),
153144
std::numeric_limits<int>::lowest());
154-
155-
#ifdef SPIRV_1_3
156145
test<class KernelName_zMyjxUrBgeUGoxmDwhvJ>(
157146
q, input, output, ext::oneapi::multiplies<int>(), 1);
158147
test<class KernelName_SljjtroxNRaAXoVnT>(q, input, output,
@@ -161,7 +150,6 @@ int main() {
161150
ext::oneapi::bit_xor<int>(), 0);
162151
test<class KernelName_xGnAnMYHvqekCk>(q, input, output,
163152
ext::oneapi::bit_and<int>(), ~0);
164-
#endif // SPIRV_1_3
165153

166154
std::cout << "Test passed." << std::endl;
167155
}

‎SYCL/GroupAlgorithm/inclusive_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>
@@ -44,8 +42,9 @@ void test(queue q, InputContainer input, OutputContainer output,
4442
typedef typename InputContainer::value_type InputT;
4543
typedef typename OutputContainer::value_type OutputT;
4644
typedef class inclusive_scan_kernel<SpecializationKernelName, 0> kernel_name0;
47-
constexpr size_t N = input.size();
45+
constexpr size_t N = input.size(); // 128 or 12
4846
constexpr size_t G = 64;
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::inclusive_scan(input.begin(), input.begin() + G, expected.begin(),
68-
binary_op, identity);
69-
assert(std::equal(output.begin(), output.begin() + G, expected.begin()));
66+
emu::inclusive_scan(input.begin(), input.begin() + confirmRange,
67+
expected.begin(), binary_op, identity);
68+
assert(std::equal(output.begin(), output.begin() + confirmRange,
69+
expected.begin()));
7070

7171
typedef class inclusive_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::inclusive_scan(input.begin(), input.begin() + G, expected.begin(),
92-
binary_op, init);
93-
assert(std::equal(output.begin(), output.begin() + G, expected.begin()));
91+
emu::inclusive_scan(input.begin(), input.begin() + confirmRange,
92+
expected.begin(), binary_op, init);
93+
assert(std::equal(output.begin(), output.begin() + confirmRange,
94+
expected.begin()));
9495

9596
typedef class inclusive_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_inclusive_scan and

‎SYCL/GroupAlgorithm/reduce.cpp

Lines changed: 0 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -4,15 +4,6 @@
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
15-
167
#include "support.h"
178
#include <CL/sycl.hpp>
189
#include <algorithm>
@@ -87,8 +78,6 @@ int main() {
8778
std::numeric_limits<int>::max());
8879
test<class KernelNameMaximumI>(q, input, output, ext::oneapi::maximum<int>(),
8980
std::numeric_limits<int>::lowest());
90-
91-
#ifdef SPIRV_1_3
9281
test<class KernelName_WonwuUVPUPOTKRKIBtT>(q, input, output,
9382
ext::oneapi::multiplies<int>(), 1);
9483
test<class KernelName_qYBaJDZTMGkdIwD>(q, input, output,
@@ -97,7 +86,6 @@ int main() {
9786
0);
9887
test<class KernelName_uFhJnxSVhNAiFPTG>(q, input, output,
9988
ext::oneapi::bit_and<int>(), ~0);
100-
#endif // SPIRV_1_3
10189

10290
std::cout << "Test passed." << std::endl;
10391
}

‎SYCL/GroupAlgorithm/reduce_sycl2020.cpp

Lines changed: 0 additions & 4 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>
@@ -82,13 +80,11 @@ int main() {
8280
test<class KernelNameMaximumI>(q, input, output, sycl::maximum<int>(),
8381
std::numeric_limits<int>::lowest());
8482

85-
#ifdef SPIRV_1_3
8683
test<class KernelNameMultipliesI>(q, input, output, sycl::multiplies<int>(),
8784
1);
8885
test<class KernelNameBitOrI>(q, input, output, sycl::bit_or<int>(), 0);
8986
test<class KernelNameBitXorI>(q, input, output, sycl::bit_xor<int>(), 0);
9087
test<class KernelNameBitAndI>(q, input, output, sycl::bit_and<int>(), ~0);
91-
#endif // SPIRV_1_3
9288

9389
// as part of SYCL_EXT_ONEAPI_COMPLEX_ALGORITHMS (
9490
// https://github.com/intel/llvm/pull/5108/ ) joint_reduce and

0 commit comments

Comments
 (0)