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

[via cfe-commits]

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(