1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
|
--- a/tests/regression/CMakeLists.txt
+++ b/tests/regression/CMakeLists.txt
@@ -49,6 +49,7 @@ set(PROGRAMS_TO_BUILD test_barrier_betwe
test_flatten_barrier_subs test_alignment_with_dynamic_wg
test_alignment_with_dynamic_wg2 test_alignment_with_dynamic_wg3
test_issue_893 test_issue_1435 test_builtin_args test_issue_1390
+ test_issue_1608
test_workitem_func_outside_kernel
)
@@ -85,6 +86,8 @@ add_test_pocl(NAME "regression/test_issu
add_test_pocl(NAME "regression/test_issue_1435" COMMAND "test_issue_1435")
+add_test_pocl(NAME "regression/test_issue_1608" COMMAND "test_issue_1608")
+
add_test_pocl(NAME "regression/test_workitem_func_outside_kernel" COMMAND "test_workitem_func_outside_kernel")
if(OPENCL_HEADER_VERSION GREATER 299)
--- /dev/null
+++ b/tests/regression/test_issue_1608.cpp
@@ -0,0 +1,63 @@
+// https://github.com/pocl/pocl/issues/1608
+// triggers on some LLVM versions
+// "LLVM ERROR: Instruction Combining did not reach a fixpoint after 1 iterations"
+
+#define CL_HPP_ENABLE_EXCEPTIONS
+#define CL_HPP_MINIMUM_OPENCL_VERSION 120
+#define CL_HPP_TARGET_OPENCL_VERSION 120
+#include <CL/opencl.hpp>
+#include <iostream>
+
+const char *SOURCE = R"RAW(
+float fn1(float *b, int c, int d) {
+ float e;
+ if (c * d)
+ e = b[0];
+ return e;
+}
+void fn2(int g, int h) {
+ float B[8], C[8], *p = 0;
+ int i = get_local_id(0);
+ for (int k = 0; k < g; ++k) {
+ int l = get_local_id(0) * get_local_id(1) * i * i * i;
+ p[l / 8 + l % 8] = B[0];
+ }
+ float f = fn1(0, i / h, i % h);
+ for (int m = 0; m < 8; ++m)
+ for (int n = 0; n < 2; ++n)
+ C[m] = mad(B[m], f, C[m]);
+}
+__kernel void krnl(int g) {
+ fn2(g, 1);
+}
+)RAW";
+
+int main(int argc, char *argv[]) {
+ int err = EXIT_SUCCESS;
+
+ try {
+ cl::Program program(SOURCE);
+ try {
+ program.build();
+ } catch (cl::BuildError &e) {
+ std::cout << "FAIL with BUILD ERROR = " << e.err() << " " << e.what() << std::endl;
+ for (auto &bl : e.getBuildLog())
+ std::cout << std::get<1>(bl);
+ return EXIT_FAILURE;
+ }
+ // This triggers compilation of dynamic WG binaries.
+ cl::Program::Binaries binaries{};
+ err = program.getInfo<>(CL_PROGRAM_BINARIES, &binaries);
+ } catch (cl::Error &e) {
+ std::cout << "FAIL with OpenCL error = " << e.err() << " " << e.what() << std::endl;
+ return EXIT_FAILURE;
+ }
+
+ if (err == CL_SUCCESS) {
+ printf("OK\n");
+ return EXIT_SUCCESS;
+ } else {
+ printf("FAIL\n");
+ return EXIT_FAILURE;
+ }
+}
|
