[libclc] [libclc] Add initial support for libclc execution conformance tests (PR #214072)
Joseph Huber via cfe-commits
cfe-commits at lists.llvm.org
Tue Aug 4 14:01:25 PDT 2026
https://github.com/jhuber6 created https://github.com/llvm/llvm-project/pull/214072
Summary:
This uses the `llvm-gpu-loader` tool to invoke OpenCL test kernels. This
routes OpenCL through the standard compute runtimes, which aren't
exactly 1-to-1 with OpenCL, but it's close enough. Obviously, the best
solution would be to have a real OpenCL host library that can launch
these, but this is the best we can get in-tree and is very sipmle, write
kernel, test kernel.
The goal is to have a lot more tests run here, but this just gets the
basic infrastructure in place. We won't be able to do full conformance
like external suites, like exhaustive math, but I think its' a step in
the right direction. It looks like this on my side:
```console
$ ninja check-libclc-amdgpu-amd-amdhsa-llvm
```
>From 4bd9db98bcc37e2ef339b6a4608b87553b2db9df Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Tue, 4 Aug 2026 15:39:19 -0500
Subject: [PATCH] [libclc] Add initial support for libclc execution conformance
tests
Summary:
This uses the `llvm-gpu-loader` tool to invoke OpenCL test kernels. This
routes OpenCL through the standard compute runtimes, which aren't
exactly 1-to-1 with OpenCL, but it's close enough. Obviously, the best
solution would be to have a real OpenCL host library that can launch
these, but this is the best we can get in-tree and is very sipmle, write
kernel, test kernel.
The goal is to have a lot more tests run here, but this just gets the
basic infrastructure in place. We won't be able to do full conformance
like external suites, like exhaustive math, but I think its' a step in
the right direction. It looks like this on my side:
```console
$ ninja check-libclc-amdgpu-amd-amdhsa-llvm
```
---
libclc/test/CMakeLists.txt | 27 ++++++++++++++++++++
libclc/test/conformance/lit.local.cfg | 2 ++
libclc/test/conformance/work_group_reduce.cl | 9 +++++++
libclc/test/lit.cfg.py | 23 ++++++++++++++++-
libclc/test/lit.site.cfg.py.in | 2 ++
5 files changed, 62 insertions(+), 1 deletion(-)
create mode 100644 libclc/test/conformance/lit.local.cfg
create mode 100644 libclc/test/conformance/work_group_reduce.cl
diff --git a/libclc/test/CMakeLists.txt b/libclc/test/CMakeLists.txt
index 0cbbc2838d624..609b3d50d1b2f 100644
--- a/libclc/test/CMakeLists.txt
+++ b/libclc/test/CMakeLists.txt
@@ -6,6 +6,33 @@ else()
set(LIBCLC_TARGET_CPU "")
endif()
+include(CheckCXXCompilerFlag)
+
+set(LIBCLC_TEST_ARCHITECTURE "" CACHE STRING
+ "Architecture for libclc execution tests; 'native' or a specific name")
+
+set(LIBCLC_OFFLOAD_LIBDIR "")
+
+if(NOT LIBCLC_TEST_ARCHITECTURE)
+ check_cxx_compiler_flag("-march=native" LIBCLC_HAS_NATIVE_DEVICE)
+ if(LIBCLC_HAS_NATIVE_DEVICE)
+ set(LIBCLC_TEST_ARCHITECTURE "native")
+ endif()
+endif()
+
+if(LIBCLC_TEST_ARCHITECTURE)
+ find_library(LIBCLC_OFFLOAD_LIB LLVMOffload
+ PATHS "${LLVM_LIBRARY_DIR}"
+ PATH_SUFFIXES "${LLVM_HOST_TRIPLE}"
+ NO_DEFAULT_PATH)
+ if(LIBCLC_OFFLOAD_LIB)
+ get_filename_component(LIBCLC_OFFLOAD_LIBDIR "${LIBCLC_OFFLOAD_LIB}" DIRECTORY)
+ else()
+ set(LIBCLC_TEST_ARCHITECTURE "")
+ message(STATUS "libclc execution tests disabled: offload runtime not found")
+ endif()
+endif()
+
configure_lit_site_cfg(
${CMAKE_CURRENT_SOURCE_DIR}/lit.site.cfg.py.in
${CMAKE_CURRENT_BINARY_DIR}/lit.site.cfg.py
diff --git a/libclc/test/conformance/lit.local.cfg b/libclc/test/conformance/lit.local.cfg
new file mode 100644
index 0000000000000..f941ad3fc9fb1
--- /dev/null
+++ b/libclc/test/conformance/lit.local.cfg
@@ -0,0 +1,2 @@
+if "libclc-native-run" not in config.available_features:
+ config.unsupported = True
diff --git a/libclc/test/conformance/work_group_reduce.cl b/libclc/test/conformance/work_group_reduce.cl
new file mode 100644
index 0000000000000..43130d802652c
--- /dev/null
+++ b/libclc/test/conformance/work_group_reduce.cl
@@ -0,0 +1,9 @@
+// RUN: %libclc-compile-and-run --threads-x 64 %t
+
+__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");
+}
diff --git a/libclc/test/lit.cfg.py b/libclc/test/lit.cfg.py
index be5cad7f7494e..fbdcdabedcd32 100644
--- a/libclc/test/lit.cfg.py
+++ b/libclc/test/lit.cfg.py
@@ -59,6 +59,9 @@ def calculate_arch_features(arch_string):
llvm_config.add_tool_substitutions(["llvm-nm"], config.llvm_tools_dir)
+target_arch = config.libclc_target_arch.lower()
+check_prefix = "AMDGCN" if target_arch == "amdgpu" else target_arch.upper()
+
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 ""
@@ -69,10 +72,28 @@ def calculate_arch_features(arch_string):
("%target", config.libclc_target),
("%cpu", config.libclc_target_cpu),
("%libclc_lib", libclc_lib),
- ("%check_prefix", config.libclc_target_arch.upper()),
+ ("%check_prefix", check_prefix),
]
)
+test_arch = getattr(config, "libclc_test_arch", "")
+offload_libdir = getattr(config, "libclc_offload_libdir", "")
+
+if test_arch and offload_libdir and os.path.isfile(path):
+ compile_cmd = (
+ f"{os.path.join(config.llvm_tools_dir, '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}"))
+ config.substitutions.append(("%libclc-compile", compile_cmd))
+ config.substitutions.append(("%libclc-run", run_cmd))
+ config.available_features.add("libclc-native-run")
+
# Propagate PATH from environment
if "PATH" in os.environ:
config.environment["PATH"] = os.path.pathsep.join(
diff --git a/libclc/test/lit.site.cfg.py.in b/libclc/test/lit.site.cfg.py.in
index 6b8b5456b2174..a401bcf031169 100644
--- a/libclc/test/lit.site.cfg.py.in
+++ b/libclc/test/lit.site.cfg.py.in
@@ -10,6 +10,8 @@ config.libclc_standalone_build = "@LIBCLC_STANDALONE_BUILD@"
config.libclc_target = "@LIBCLC_TARGET@"
config.libclc_target_arch = "@LIBCLC_TARGET_ARCH@"
config.libclc_target_cpu = "@LIBCLC_TARGET_CPU@"
+config.libclc_test_arch = "@LIBCLC_TEST_ARCHITECTURE@"
+config.libclc_offload_libdir = "@LIBCLC_OFFLOAD_LIBDIR@"
import lit.llvm
lit.llvm.initialize(lit_config, config)
More information about the cfe-commits
mailing list