1- // XFAIL: cuda || opencl
1+ // XFAIL: cuda
2+ // The negative test fails on CUDA. It's not clear whether the CUDA backend
3+ // respects the reqd_work_group_size attribute.
4+
25// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
36// RUN: %CPU_RUN_PLACEHOLDER %t.out
47// RUN: %GPU_RUN_PLACEHOLDER %t.out
@@ -20,59 +23,26 @@ int main() {
2023 queue Q (AsyncHandler);
2124 device D (Q.get_device ());
2225
23- string_class DeviceVendorName = D.get_info <info::device::vendor>();
24- auto DeviceType = D.get_info <info::device::device_type>();
25-
26- // parallel_for, (16, 16, 16) global, (8, 8, 8) local, reqd_wg_size(4, 4, 4)
27- // -> fail
28- try {
29- Q.submit ([&](handler &CGH) {
30- CGH.parallel_for <class ReqdWGSizeNegativeA >(
31- nd_range<3 >(range<3 >(16 , 16 , 16 ), range<3 >(8 , 8 , 8 )),
32- [=](nd_item<3 >) [[intel::reqd_work_group_size (4 , 4 , 4 )]]{
26+ bool IsOpenCL = (D.get_platform ().get_backend () == backend::opencl);
3327
34- });
35- });
36- Q.wait_and_throw ();
37- std::cerr << " Test case ReqdWGSizeNegativeA failed: no exception has been "
38- " thrown\n "
39- return 1 ; // We shouldn't be here, exception is expected
40- } catch (nd_range_error &E) {
41- if (string_class (E.what ()).find (
42- " Specified local size doesn't match the required work-group size "
43- " specified in the program source"
44- std::cerr
45- << " Test case ReqdWGSizeNegativeA failed 1: unexpected exception: "
46- << E.what () << std::endl;
47- return 1 ;
48- }
49- } catch (runtime_error &E) {
50- std::cerr << " Test case ReqdWGSizeNegativeA failed 2: unexpected exception: "
51- << E.what () << std::endl;
52- return 1 ;
53- } catch (...) {
54- std::cerr << " Test case ReqdWGSizeNegativeA failed: something unexpected "
55- " has been caught"
56- << std::endl;
57- return 1 ;
58- }
59-
60- // Positive test-cases that should pass on any underlying OpenCL runtime
28+ // Positive test case: Specify local size that matches required size.
6129 // parallel_for, (8, 8, 8) global, (4, 4, 4) local, reqd_wg_size(4, 4, 4) ->
6230 // pass
6331 try {
6432 Q.submit ([&](handler &CGH) {
6533 CGH.parallel_for <class ReqdWGSizePositiveA >(
66- nd_range<3 >(range<3 >(8 , 8 , 8 ), range<3 >(4 , 4 , 4 )),
67- [= ](nd_item<3 >) [[intel::reqd_work_group_size (4 , 4 , 4 )]]{});
34+ nd_range<3 >(range<3 >(8 , 8 , 8 ), range<3 >(4 , 4 , 4 )), [=
35+ ](nd_item<3 >) [[intel::reqd_work_group_size (4 , 4 , 4 )]]{});
6836 });
6937 Q.wait_and_throw ();
7038 } catch (nd_range_error &E) {
71- std::cerr << " Test case ReqdWGSizePositiveA failed: unexpected exception: "
39+ std::cerr << " Test case ReqdWGSizePositiveA failed: unexpected "
40+ " nd_range_error exception: "
7241 << E.what () << std::endl;
7342 return 1 ;
7443 } catch (runtime_error &E) {
75- std::cerr << " Test case ReqdWGSizePositiveA failed: unexpected exception: "
44+ std::cerr << " Test case ReqdWGSizePositiveA failed: unexpected "
45+ " runtime_error exception: "
7646 << E.what () << std::endl;
7747 return 1 ;
7848 } catch (...) {
@@ -82,24 +52,68 @@ int main() {
8252 return 1 ;
8353 }
8454
55+ // Kernel that has a required WG size, but no local size is specified.
56+ //
57+ // TODO: This fails on OpenCL and should be investigated.
58+ if (!IsOpenCL) {
59+ try {
60+ Q.submit ([&](handler &CGH) {
61+ CGH.parallel_for <class ReqdWGSizeNoLocalPositive >(
62+ range<3 >(16 , 16 , 16 ), [=
63+ ](item<3 >) [[intel::reqd_work_group_size (4 , 4 , 4 )]]{});
64+ });
65+ Q.wait_and_throw ();
66+ } catch (nd_range_error &E) {
67+ std::cerr << " Test case ReqdWGSizeNoLocalPositive failed: unexpected "
68+ " nd_range_error exception: "
69+ << E.what () << std::endl;
70+ return 1 ;
71+ } catch (runtime_error &E) {
72+ std::cerr
73+ << " Test case ReqdWGSizeNoLocalPositive: unexpected runtime_error "
74+ " exception: "
75+ << E.what () << std::endl;
76+ return 1 ;
77+ } catch (...) {
78+ std::cerr << " Test case ReqdWGSizeNoLocalPositive failed: something "
79+ " unexpected has been caught"
80+ << std::endl;
81+ return 1 ;
82+ }
83+ }
84+
85+ // Negative test case: Specify local size that does not match required size.
86+ // parallel_for, (16, 16, 16) global, (8, 8, 8) local, reqd_wg_size(4, 4, 4)
87+ // -> fail
8588 try {
8689 Q.submit ([&](handler &CGH) {
87- CGH.parallel_for <class ReqdWGSizePositiveB >(
88- range<3 >(16 , 16 , 16 ), [=](item<3 >) [[intel::reqd_work_group_size (4 , 4 , 4 )]]{});
90+ CGH.parallel_for <class ReqdWGSizeNegativeA >(
91+ nd_range<3 >(range<3 >(16 , 16 , 16 ), range<3 >(8 , 8 , 8 )), [=
92+ ](nd_item<3 >) [[intel::reqd_work_group_size (4 , 4 , 4 )]]{
93+
94+ });
8995 });
9096 Q.wait_and_throw ();
91-
97+ std::cerr << " Test case ReqdWGSizeNegativeA failed: no exception has been "
98+ " thrown\n "
99+ return 1 ; // We shouldn't be here, exception is expected
92100 } catch (nd_range_error &E) {
93- std::cerr << " Test case ReqdWGSizePositiveB failed 1: unexpected exception: "
94- << E.what () << std::endl;
95- return 1 ;
101+ if (string_class (E.what ()).find (
102+ " Specified local size doesn't match the required work-group size "
103+ " specified in the program source"
104+ std::cerr
105+ << " Test case ReqdWGSizeNegativeA failed: unexpected nd_range_error "
106+ " exception: "
107+ << E.what () << std::endl;
108+ return 1 ;
109+ }
96110 } catch (runtime_error &E) {
97- std::cerr
98- << " Test case ReqdWGSizePositiveB failed 2: unexpected exception: "
99- << E.what () << std::endl;
111+ std::cerr << " Test case ReqdWGSizeNegativeA failed: unexpected "
112+ " nd_range_error exception: "
113+ << E.what () << std::endl;
100114 return 1 ;
101115 } catch (...) {
102- std::cerr << " Test case ReqdWGSizePositiveB failed: something unexpected "
116+ std::cerr << " Test case ReqdWGSizeNegativeA failed: something unexpected "
103117 " has been caught"
104118 << std::endl;
105119 return 1 ;
0 commit comments