From 8976b387977511298cab0146af88eaebfc6615a6 Mon Sep 17 00:00:00 2001 From: Debian OpenCL Maintainers Date: Thu, 9 Jan 2025 12:44:04 +0100 Subject: [PATCH] issue_1608 Gbp-Pq: Name issue_1608.patch --- tests/regression/CMakeLists.txt | 3 ++ tests/regression/test_issue_1608.cpp | 63 ++++++++++++++++++++++++++++ 2 files changed, 66 insertions(+) create mode 100644 tests/regression/test_issue_1608.cpp diff --git a/tests/regression/CMakeLists.txt b/tests/regression/CMakeLists.txt index d51cb2e..23e202c 100644 --- a/tests/regression/CMakeLists.txt +++ b/tests/regression/CMakeLists.txt @@ -49,6 +49,7 @@ set(PROGRAMS_TO_BUILD test_barrier_between_for_loops test_early_return 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_issue_757" COMMAND "test_issue_757") 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) diff --git a/tests/regression/test_issue_1608.cpp b/tests/regression/test_issue_1608.cpp new file mode 100644 index 0000000..bbd4fd8 --- /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 +#include + +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; + } +} -- 2.30.2