https://github.com/jhuber6 created https://github.com/llvm/llvm-project/pull/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. >From 44f573303c66efd3817eaee7ee5f3bcc7e8335f6 Mon Sep 17 00:00:00 2001 From: Joseph Huber <[email protected]> Date: Wed, 5 Aug 2026 07:34:06 -0500 Subject: [PATCH] [libclc] Improve libclc device tests with feature checks 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. --- libclc/test/conformance/work_group_reduce.cl | 3 ++- libclc/test/lit.cfg.py | 25 ++++++++++++++++---- 2 files changed, 23 insertions(+), 5 deletions(-) diff --git a/libclc/test/conformance/work_group_reduce.cl b/libclc/test/conformance/work_group_reduce.cl index 43130d802652c..cd1f688884649 100644 --- a/libclc/test/conformance/work_group_reduce.cl +++ b/libclc/test/conformance/work_group_reduce.cl @@ -1,4 +1,5 @@ -// 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 __kernel void test(void) { uint lid = get_local_id(0); diff --git a/libclc/test/lit.cfg.py b/libclc/test/lit.cfg.py index 237e2418075bf..bc4f52e88a6b1 100644 --- a/libclc/test/lit.cfg.py +++ b/libclc/test/lit.cfg.py @@ -3,6 +3,7 @@ """ import os +import subprocess import lit.formats @@ -62,6 +63,8 @@ def calculate_arch_features(arch_string): target_arch = config.libclc_target_arch.lower() check_prefix = "AMDGCN" if target_arch == "amdgpu" else target_arch.upper() +config.available_features.add(target_arch) + is_standalone = config.libclc_standalone_build.lower() == "true" path = os.path.join(config.libclc_library_dir, config.libclc_target, "libclc.bc") libclc_lib = f"--libclc-lib=:{path}" if is_standalone else "" @@ -80,22 +83,36 @@ 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
