Skip to content

Commit 17eb598

Browse files
fadeevalfadeeval
authored
[SYCL] Add negative tests for inline asm (#2406)
1 parent 67acf81 commit 17eb598

11 files changed

+444
-19
lines changed

‎sycl/test/inline-asm/include/asmhelper.h

Lines changed: 44 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -84,35 +84,60 @@ bool isInlineASMSupported(sycl::device Device) {
8484
return true;
8585
}
8686

87-
/// checks if device suppots inline asm feature and launches a test
88-
///
89-
/// \returns false if test wasn't launched (i.e.was skipped) and true otherwise
87+
auto exception_handler = [](sycl::exception_list exceptions) {
88+
for (std::exception_ptr const &e : exceptions) {
89+
try {
90+
std::rethrow_exception(e);
91+
} catch(sycl::exception const &e) {
92+
std::cout << "Caught asynchronous SYCL exception:\n"
93+
<< e.what() << std::endl;
94+
}
95+
}
96+
};
97+
9098
template <typename F>
91-
bool launchInlineASMTest(F &f, bool requires_particular_sg_size = true) {
92-
try {
93-
cl::sycl::queue deviceQueue(cl::sycl::gpu_selector{});
94-
cl::sycl::device device = deviceQueue.get_device();
99+
bool launchInlineASMTestImpl(F &f, bool requires_particular_sg_size = true) {
100+
cl::sycl::queue deviceQueue(cl::sycl::gpu_selector{}, exception_handler);
101+
cl::sycl::device device = deviceQueue.get_device();
95102

96103
#if defined(INLINE_ASM)
97-
if (!isInlineASMSupported(device)) {
98-
std::cout << "Skipping test\n";
99-
return false;
100-
}
104+
if (!isInlineASMSupported(device)) {
105+
std::cout << "Skipping test\n";
106+
return false;
107+
}
101108
#endif
102109

103-
if (requires_particular_sg_size && !device.has_extension("cl_intel_required_subgroup_size")) {
104-
std::cout << "Skipping test\n";
105-
return false;
106-
}
107-
108-
deviceQueue.submit(f).wait_and_throw();
109-
} catch (cl::sycl::exception &e) {
110-
std::cerr << "Caught exception: " << e.what() << std::endl;
110+
if (requires_particular_sg_size &&
111+
!device.has_extension("cl_intel_required_subgroup_size")) {
112+
std::cout << "Skipping test\n";
113+
return false;
111114
}
112115

116+
deviceQueue.submit(f).wait_and_throw();
117+
113118
return true;
114119
}
115120

121+
/// checks if device suppots inline asm feature and launches a test
122+
///
123+
/// \returns false if test wasn't launched (i.e.was skipped) and true otherwise
124+
template <typename F>
125+
bool launchInlineASMTest(F &f, bool requires_particular_sg_size = true,
126+
bool need_to_throw_exception = false) {
127+
if (need_to_throw_exception) {
128+
return launchInlineASMTestImpl(f, requires_particular_sg_size);
129+
} else {
130+
bool result = false;
131+
try {
132+
result = launchInlineASMTestImpl(f, requires_particular_sg_size);
133+
} catch (cl::sycl::exception &e) {
134+
std::cerr << "Caught exception: " << e.what() << std::endl;
135+
}
136+
137+
return result;
138+
}
139+
}
140+
116141
template <typename T>
117142
bool verify_all_the_same(const std::vector<T> &input, T reference_value) {
118143
for (int i = 0; i < input.size(); ++i)

‎sycl/test/inline-asm/negative_tests/asm_bad_opcode.cpp

Lines changed: 40 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,40 @@
1+
// UNSUPPORTED: cuda
2+
// REQUIRES: gpu,linux
3+
// RUN: %clangxx -fsycl %s -o %t.out
4+
// TODO: enable the line below once we update NEO driver in our CI
5+
// RUNx: %t.out
6+
7+
#include "../include/asmhelper.h"
8+
#include <CL/sycl.hpp>
9+
10+
struct KernelFunctor {
11+
KernelFunctor() {}
12+
13+
void operator()(cl::sycl::handler &cgh) {
14+
cgh.parallel_for<KernelFunctor>(
15+
cl::sycl::range<1>{16}, [=](cl::sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
16+
#if defined(__SYCL_DEVICE_ONLY__)
17+
asm volatile(".decl tmp1 v_type=G type=d num_elts=16 align=GRF\n"
18+
".decl tmp2 v_type=G type=d num_elts=16 align=GRF\n"
19+
"movi (M1_NM, 8) tmp1(0,1)<1> tmp2(0,0)\n");
20+
#endif
21+
});
22+
}
23+
};
24+
25+
int main() {
26+
KernelFunctor f;
27+
try {
28+
launchInlineASMTest(f, /* sg size */ true,
29+
/* exception is expected */ true);
30+
} catch (const cl::sycl::compile_program_error &e) {
31+
std::string what = e.what();
32+
// TODO: check for precise exception class and message once they are known
33+
// (pending driver update)
34+
if (what.find("syntax error") != std::string::npos) {
35+
return 0;
36+
}
37+
}
38+
std::cout << "Expected an exception about syntax error" << std::endl;
39+
return 1;
40+
}

‎sycl/test/inline-asm/negative_tests/asm_bad_operand_syntax.cpp

Lines changed: 40 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,40 @@
1+
// UNSUPPORTED: cuda
2+
// REQUIRES: gpu,linux
3+
// RUN: %clangxx -fsycl %s -o %t.out
4+
// TODO: enable the line below once we update NEO driver in our CI
5+
// RUNx: %t.out
6+
7+
#include "../include/asmhelper.h"
8+
#include <CL/sycl.hpp>
9+
10+
struct KernelFunctor {
11+
KernelFunctor() {}
12+
13+
void operator()(cl::sycl::handler &cgh) {
14+
cgh.parallel_for<KernelFunctor>(
15+
cl::sycl::range<1>{16}, [=](cl::sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
16+
#if defined(__SYCL_DEVICE_ONLY__)
17+
asm volatile(".decl tmp1 v_type=G type=d num_elts=16 align=GRF\n"
18+
".decl tmp2 v_type=G type=d num_elts=16 align=GRF\n"
19+
"mov (M1_NM, 8) tmp1(0,1)<1>:f tmp2(0,0)<1;1,0>\n");
20+
#endif
21+
});
22+
}
23+
};
24+
25+
int main() {
26+
KernelFunctor f;
27+
try {
28+
launchInlineASMTest(f, /* sg size */ true,
29+
/* exception is expected */ true);
30+
} catch (const cl::sycl::compile_program_error &e) {
31+
std::string what = e.what();
32+
// TODO: check for precise exception class and message once they are known
33+
// (pending driver update)
34+
if (what.find("syntax error") != std::string::npos) {
35+
return 0;
36+
}
37+
}
38+
std::cout << "Expected an exception about syntax error" << std::endl;
39+
return 1;
40+
}

‎sycl/test/inline-asm/negative_tests/asm_duplicate_label.cpp

Lines changed: 40 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,40 @@
1+
// UNSUPPORTED: cuda
2+
// REQUIRES: gpu,linux
3+
// RUN: %clangxx -fsycl %s -o %t.out
4+
// TODO: enable the line below once we update NEO driver in our CI
5+
// RUNx: %t.out
6+
7+
#include "../include/asmhelper.h"
8+
#include <CL/sycl.hpp>
9+
10+
struct KernelFunctor {
11+
KernelFunctor() {}
12+
13+
void operator()(cl::sycl::handler &cgh) {
14+
cgh.parallel_for<KernelFunctor>(
15+
cl::sycl::range<1>{16}, [=](cl::sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
16+
#if defined(__SYCL_DEVICE_ONLY__)
17+
asm volatile(".decl tmp1 v_type=G type=d num_elts=16 align=GRF\n"
18+
".decl tmp2 v_type=G type=d num_elts=16 align=GRF\n"
19+
"check_label0:\ncheck_label0:\n");
20+
#endif
21+
});
22+
}
23+
};
24+
25+
int main() {
26+
KernelFunctor f;
27+
try {
28+
launchInlineASMTest(f, /* sg size */ true,
29+
/* exception is expected */ true);
30+
} catch (const cl::sycl::exception &e) {
31+
std::string what = e.what();
32+
// TODO: check for precise exception class and message once they are known
33+
// (pending driver update)
34+
if (what.find("OpenCL API failed") != std::string::npos) {
35+
return 0;
36+
}
37+
}
38+
std::cout << "Expected an exception about syntax error" << std::endl;
39+
return 1;
40+
}

‎sycl/test/inline-asm/negative_tests/asm_illegal_exec_size.cpp

Lines changed: 40 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,40 @@
1+
// UNSUPPORTED: cuda
2+
// REQUIRES: gpu,linux
3+
// RUN: %clangxx -fsycl %s -o %t.out
4+
// TODO: enable the line below once we update NEO driver in our CI
5+
// RUNx: %t.out
6+
7+
#include "../include/asmhelper.h"
8+
#include <CL/sycl.hpp>
9+
10+
struct KernelFunctor {
11+
KernelFunctor() {}
12+
13+
void operator()(cl::sycl::handler &cgh) {
14+
cgh.parallel_for<KernelFunctor>(
15+
cl::sycl::range<1>{16}, [=](cl::sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
16+
#if defined(__SYCL_DEVICE_ONLY__)
17+
asm volatile(".decl tmp1 v_type=G type=d num_elts=16 align=GRF\n"
18+
".decl tmp2 v_type=G type=d num_elts=16 align=GRF\n"
19+
"mov (M1_NM, 6) tmp1(0,1)<1> tmp2(0,0)<1;1,0>\n");
20+
#endif
21+
});
22+
}
23+
};
24+
25+
int main() {
26+
KernelFunctor f;
27+
try {
28+
launchInlineASMTest(f, /* sg size */ true,
29+
/* exception is expected */ true);
30+
} catch (const cl::sycl::compile_program_error &e) {
31+
std::string what = e.what();
32+
// TODO: check for precise exception class and message once they are known
33+
// (pending driver update)
34+
if (what.find("invalid execution size") != std::string::npos) {
35+
return 0;
36+
}
37+
}
38+
std::cout << "Expected an exception about syntax error" << std::endl;
39+
return 1;
40+
}

‎sycl/test/inline-asm/negative_tests/asm_missing_label.cpp

Lines changed: 40 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,40 @@
1+
// UNSUPPORTED: cuda
2+
// REQUIRES: gpu,linux
3+
// RUN: %clangxx -fsycl %s -o %t.out
4+
// TODO: enable the line below once we update NEO driver in our CI
5+
// RUNx: %t.out
6+
7+
#include "../include/asmhelper.h"
8+
#include <CL/sycl.hpp>
9+
10+
struct KernelFunctor {
11+
KernelFunctor() {}
12+
13+
void operator()(cl::sycl::handler &cgh) {
14+
cgh.parallel_for<KernelFunctor>(
15+
cl::sycl::range<1>{16}, [=](cl::sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
16+
#if defined(__SYCL_DEVICE_ONLY__)
17+
asm volatile(".decl tmp1 v_type=G type=d num_elts=16 align=GRF\n"
18+
".decl tmp2 v_type=G type=d num_elts=16 align=GRF\n"
19+
"goto (M1, 16) check_label0\n");
20+
#endif
21+
});
22+
}
23+
};
24+
25+
int main() {
26+
KernelFunctor f;
27+
try {
28+
launchInlineASMTest(f, /* sg size */ true,
29+
/* exception is expected */ true);
30+
} catch (const cl::sycl::compile_program_error &e) {
31+
std::string what = e.what();
32+
// TODO: check for precise exception class and message once they are known
33+
// (pending driver update)
34+
if (what.find("OpenCL API failed") != std::string::npos) {
35+
return 0;
36+
}
37+
}
38+
std::cout << "Expected an exception about syntax error" << std::endl;
39+
return 1;
40+
}

‎sycl/test/inline-asm/negative_tests/asm_missing_region.cpp

Lines changed: 40 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,40 @@
1+
// UNSUPPORTED: cuda
2+
// REQUIRES: gpu,linux
3+
// RUN: %clangxx -fsycl %s -o %t.out
4+
// TODO: enable the line below once we update NEO driver in our CI
5+
// RUNx: %t.out
6+
7+
#include "../include/asmhelper.h"
8+
#include <CL/sycl.hpp>
9+
10+
struct KernelFunctor {
11+
KernelFunctor() {}
12+
13+
void operator()(cl::sycl::handler &cgh) {
14+
cgh.parallel_for<KernelFunctor>(
15+
cl::sycl::range<1>{16}, [=](cl::sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
16+
#if defined(__SYCL_DEVICE_ONLY__)
17+
asm volatile(".decl tmp1 v_type=G type=d num_elts=16 align=GRF\n"
18+
".decl tmp2 v_type=G type=d num_elts=16 align=GRF\n"
19+
"mov (M1_NM, 8) tmp1(0,1)<1> tmp2(0,0)\n");
20+
#endif
21+
});
22+
}
23+
};
24+
25+
int main() {
26+
KernelFunctor f;
27+
try {
28+
launchInlineASMTest(f, /* sg size */ true,
29+
/* exception is expected */ true);
30+
} catch (const cl::sycl::compile_program_error &e) {
31+
std::string what = e.what();
32+
// TODO: check for precise exception class and message once they are known
33+
// (pending driver update)
34+
if (what.find("syntax error") != std::string::npos) {
35+
return 0;
36+
}
37+
}
38+
std::cout << "Expected an exception about syntax error" << std::endl;
39+
return 1;
40+
}

‎sycl/test/inline-asm/negative_tests/asm_simple.cpp

Lines changed: 40 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,40 @@
1+
// UNSUPPORTED: cuda
2+
// REQUIRES: gpu,linux
3+
// RUN: %clangxx -fsycl %s -o %t.out
4+
// TODO: enable the line below once we update NEO driver in our CI
5+
// RUNx: %t.out
6+
7+
#include "../include/asmhelper.h"
8+
#include <CL/sycl.hpp>
9+
10+
struct KernelFunctor {
11+
KernelFunctor() {}
12+
13+
void operator()(cl::sycl::handler &cgh) {
14+
cgh.parallel_for<KernelFunctor>(
15+
cl::sycl::range<1>{16}, [=](cl::sycl::id<1> wiID) [[intel::reqd_sub_group_size(8)]] {
16+
#if defined(__SYCL_DEVICE_ONLY__)
17+
asm volatile(".decl tmp1 v_type=G type=d num_elts=16 align=GRF\n"
18+
".decl tmp2 v_type=G type=d num_elts=16 align=GRF\n"
19+
"@@\n");
20+
#endif
21+
});
22+
}
23+
};
24+
25+
int main() {
26+
KernelFunctor f;
27+
try {
28+
launchInlineASMTest(f, /* sg size */ true,
29+
/* exception is expected */ true);
30+
} catch (const cl::sycl::compile_program_error &e) {
31+
std::string what = e.what();
32+
// TODO: check for precise exception class and message once they are known
33+
// (pending driver update)
34+
if (what.find("syntax error") != std::string::npos) {
35+
return 0;
36+
}
37+
}
38+
std::cout << "Expected an exception about syntax error" << std::endl;
39+
return 1;
40+
}

0 commit comments

Comments
 (0)