Author: Joseph Huber Date: 2026-08-06T09:00:25-05:00 New Revision: c90a5cbb0d3366526a5d44032b06b3e624a5249b
URL: https://github.com/llvm/llvm-project/commit/c90a5cbb0d3366526a5d44032b06b3e624a5249b DIFF: https://github.com/llvm/llvm-project/commit/c90a5cbb0d3366526a5d44032b06b3e624a5249b.diff LOG: [libclc] Improve libclc device tests with feature checks (#214217) Summary: Make sure the tests can conditionally run on the correct features. I think that it would be nice if we could print from these tests, but that would require linking against `libc` for GPUs and I'm unsure if that's worth it, so for now I'm just sticking with traps as the failure mechanism. Added: libclc/test/conformance/conformance.h Modified: libclc/test/conformance/work_group_reduce.cl libclc/test/lit.cfg.py Removed: ################################################################################ diff --git a/libclc/test/conformance/conformance.h b/libclc/test/conformance/conformance.h new file mode 100644 index 0000000000000..6c98681b29e86 --- /dev/null +++ b/libclc/test/conformance/conformance.h @@ -0,0 +1,32 @@ +//===----------------------------------------------------------------------===// +// +// Shared helpers for libclc execution conformance tests. These kernels are +// launched on-device and report failure by trapping. +// +//===----------------------------------------------------------------------===// + +#ifndef LIBCLC_TEST_CONFORMANCE_H +#define LIBCLC_TEST_CONFORMANCE_H + +// Yields an input value the compiler cannot optimize out. +#define TEST_INPUT(TYPE, VALUE) \ + ({ \ + volatile TYPE __clc_in = (VALUE); \ + __clc_in; \ + }) + +// Check the condition and submit a hardware trap on failure. +#define CHECK(COND) \ + do { \ + if (!(COND)) \ + __builtin_verbose_trap("libclc", "check failed: " #COND); \ + } while (0) + +#define CHECK_EQ(LHS, RHS) CHECK((LHS) == (RHS)) +#define CHECK_NE(LHS, RHS) CHECK((LHS) != (RHS)) +#define CHECK_LT(LHS, RHS) CHECK((LHS) < (RHS)) +#define CHECK_LE(LHS, RHS) CHECK((LHS) <= (RHS)) +#define CHECK_GT(LHS, RHS) CHECK((LHS) > (RHS)) +#define CHECK_GE(LHS, RHS) CHECK((LHS) >= (RHS)) + +#endif // LIBCLC_TEST_CONFORMANCE_H diff --git a/libclc/test/conformance/work_group_reduce.cl b/libclc/test/conformance/work_group_reduce.cl index 43130d802652c..98662d319c5cb 100644 --- a/libclc/test/conformance/work_group_reduce.cl +++ b/libclc/test/conformance/work_group_reduce.cl @@ -1,9 +1,12 @@ -// RUN: %libclc-compile-and-run --threads-x 64 %t +// REQUIRES: __opencl_c_work_group_collective_functions +// RUN: %libclc-compile-and-run --kernel test --threads-x 64 %t + +#include "conformance.h" __kernel void test(void) { uint lid = get_local_id(0); uint n = get_local_size(0); uint sum = work_group_reduce_add(lid); - if (lid == 0 && sum != n * (n - 1) / 2) - __builtin_verbose_trap("libclc", "work_group_reduce_add mismatch"); + if (lid == 0) + CHECK_EQ(sum, n * (n - 1) / 2); } diff --git a/libclc/test/lit.cfg.py b/libclc/test/lit.cfg.py index 237e2418075bf..d6c49dbfe5878 100644 --- a/libclc/test/lit.cfg.py +++ b/libclc/test/lit.cfg.py @@ -3,6 +3,7 @@ """ import os +import subprocess import lit.formats @@ -80,22 +81,47 @@ def calculate_arch_features(arch_string): offload_libdir = getattr(config, "libclc_offload_libdir", "") if test_arch and offload_libdir and os.path.isfile(path): + clang = os.path.join(config.llvm_tools_dir, "clang") + loader = os.path.join(config.llvm_tools_dir, "llvm-gpu-loader") compile_cmd = ( - f"{os.path.join(config.llvm_tools_dir, 'clang')} --target={config.libclc_target} " + f"{clang} --target={config.libclc_target} " f"-march={test_arch} -cl-std=CL3.0 -nogpulib " f"--libclc-lib=:{path} %s -o %t" ) - run_cmd = f"{os.path.join(config.llvm_tools_dir, 'llvm-gpu-loader')} --kernel test" config.environment["LD_LIBRARY_PATH"] = os.pathsep.join( [offload_libdir, config.environment.get("LD_LIBRARY_PATH", "")] ) config.substitutions.append( - ("%libclc-compile-and-run", f"{compile_cmd} && {run_cmd}") + ("%libclc-compile-and-run", f"{compile_cmd} && {loader}") ) config.substitutions.append(("%libclc-compile", compile_cmd)) - config.substitutions.append(("%libclc-run", run_cmd)) + config.substitutions.append(("%libclc-run", loader)) config.available_features.add("libclc-native-run") + # Register the OpenCL features and extensions the compiler advertises. + probe = subprocess.run( + [ + clang, + f"--target={config.libclc_target}", + f"-march={test_arch}", + "-cl-std=CL3.0", + "-x", + "cl", + "-dM", + "-E", + "-", + ], + input="", + capture_output=True, + text=True, + ) + for line in probe.stdout.splitlines(): + tokens = line.split() + if len(tokens) >= 2 and tokens[0] == "#define": + name = tokens[1] + if name.startswith("__opencl_c_") or name.startswith("cl_khr_"): + config.available_features.add(name) + # Propagate PATH from environment if "PATH" in os.environ: config.environment["PATH"] = os.path.pathsep.join( _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
