[libclc] [libclc] Improve libclc device tests with feature checks (PR #214217)

[Joseph Huber via cfe-commits]

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 <huberjn at outlook.com>
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(