Author: Joseph Huber Date: 2026-08-04T22:47:00-05:00 New Revision: d87c685d5d88fe00b0594fc238b9afb16ba6eb64
URL: https://github.com/llvm/llvm-project/commit/d87c685d5d88fe00b0594fc238b9afb16ba6eb64 DIFF: https://github.com/llvm/llvm-project/commit/d87c685d5d88fe00b0594fc238b9afb16ba6eb64.diff LOG: [libclc] Add initial support for libclc execution conformance tests (#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 ``` This requires building with the LLVM `offload/` runtime enabled for the source compiler. Added: libclc/test/conformance/lit.local.cfg libclc/test/conformance/work_group_reduce.cl Modified: libclc/test/CMakeLists.txt libclc/test/lit.cfg.py libclc/test/lit.site.cfg.py.in Removed: ################################################################################ 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..237e2418075bf 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,30 @@ 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) _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
