[libclc] Add initial LIT tests

These tests aren't very meaningful and aren't immune to false positives,
but they do get the project building when running 'check-all' and so
enable libclc testing in CI.

Author: frasercrmck

libclc/CMakeLists.txt

@@ -366,7 +366,10 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} )
         COMMAND ${LLVM_SPIRV} ${spvflags} -o ${spv_suffix} ${builtins_link_lib}
         DEPENDS ${builtins_link_lib}
       )
-      add_custom_target( "prepare-${spv_suffix}" ALL DEPENDS "${spv_suffix}" )
+      add_custom_target( prepare-${arch_suffix} ALL DEPENDS ${spv_suffix} )
+      set_target_properties( prepare-${arch_suffix}
+        PROPERTIES TARGET_FILE ${spv_suffix}
+      )
       install( FILES ${CMAKE_CURRENT_BINARY_DIR}/${spv_suffix}
          DESTINATION "${CMAKE_INSTALL_DATADIR}/clc" )
     else()
@@ -392,7 +395,10 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} )
       add_custom_command( OUTPUT ${obj_suffix}
         COMMAND prepare_builtins -o ${obj_suffix} ${builtins_opt_lib}
         DEPENDS ${builtins_opt_lib} prepare_builtins )
-      add_custom_target( prepare-${obj_suffix} ALL DEPENDS ${obj_suffix} )
+      add_custom_target( prepare-${arch_suffix} ALL DEPENDS ${obj_suffix} )
+      set_target_properties( prepare-${arch_suffix}
+        PROPERTIES TARGET_FILE ${obj_suffix}
+      )
 
       # nvptx-- targets don't include workitem builtins
       if( NOT clang_triple MATCHES ".*ptx.*--$" )
@@ -406,9 +412,13 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} )
         set( alias_suffix "${a}-${clang_triple}.bc" )
         add_custom_target( ${alias_suffix} ALL
           COMMAND ${CMAKE_COMMAND} -E create_symlink ${obj_suffix} ${alias_suffix}
-          DEPENDS prepare-${obj_suffix} )
+          DEPENDS prepare-${arch_suffix} )
         install( FILES ${CMAKE_CURRENT_BINARY_DIR}/${alias_suffix} DESTINATION "${CMAKE_INSTALL_DATADIR}/clc" )
       endforeach( a )
     endif()
   endforeach( d )
 endforeach( t )
+
+if( NOT LIBCLC_STANDALONE_BUILD )
+  add_subdirectory( test )
+endif()

libclc/test/CMakeLists.txt

@@ -0,0 +1,46 @@
+set(LIBCLC_TEST_SOURCE_DIR ${CMAKE_CURRENT_SOURCE_DIR})
+
+configure_lit_site_cfg(
+  ${CMAKE_CURRENT_SOURCE_DIR}/lit.site.cfg.py.in
+  ${CMAKE_CURRENT_BINARY_DIR}/lit.site.cfg.py
+  MAIN_CONFIG
+  ${CMAKE_CURRENT_SOURCE_DIR}/lit.cfg.py
+)
+
+list(APPEND LIBCLC_TEST_DEPS
+  libclc::clang
+  FileCheck count not
+)
+
+add_custom_target(check-libclc)
+
+foreach( t ${LIBCLC_TARGETS_TO_BUILD} )
+  foreach( d ${${t}_devices} )
+
+    # The tests use LLVM IR, so we can't test SPIR-V libraries here.
+    string( REPLACE "-" ";" TRIPLE  ${t} )
+    list( GET TRIPLE 0 ARCH )
+    if( ARCH STREQUAL spirv OR ARCH STREQUAL spirv64 )
+      continue()
+    endif()
+
+    get_libclc_device_info(
+      TRIPLE ${t}
+      DEVICE ${d}
+      CPU cpu
+      ARCH_SUFFIX arch_suffix
+      CLANG_TRIPLE clang_triple
+    )
+
+    add_lit_testsuite(check-libclc-${arch_suffix}
+      "Running libclc ${arch_suffix} regression tests"
+      ${CMAKE_CURRENT_BINARY_DIR}
+      PARAMS "target=${clang_triple}" cpu=${cpu}
+        "builtins=$<TARGET_PROPERTY:prepare-${arch_suffix},TARGET_FILE>"
+      DEPENDS prepare-${arch_suffix} ${LIBCLC_TEST_DEPS}
+    )
+
+    add_dependencies( check-libclc check-libclc-${arch_suffix} )
+
+  endforeach()
+endforeach()

libclc/test/add_sat.cl

@@ -1,3 +1,6 @@
+// RUN: %clang -emit-llvm -S -o - %s | FileCheck %s
+
+// CHECK: foo
 __kernel void foo(__global char *a, __global char *b, __global char *c) {
   *a = add_sat(*b, *c);
 }

libclc/test/as_type.cl

@@ -1,3 +1,6 @@
-__kernel void foo(int4 *x, float4 *y) {
+// RUN: %clang -emit-llvm -S -o - %s | FileCheck %s
+
+// CHECK: foo
+__kernel void foo(__global int4 *x, __global float4 *y) {
   *x = as_int4(*y);
 }

libclc/test/convert.cl

@@ -1,3 +1,6 @@
-__kernel void foo(int4 *x, float4 *y) {
+// RUN: %clang -emit-llvm -S -o - %s | FileCheck %s
+
+// CHECK: foo
+__kernel void foo(__global int4 *x, __global float4 *y) {
   *x = convert_int4(*y);
 }

libclc/test/cos.cl

@@ -1,3 +1,6 @@
-__kernel void foo(float4 *f) {
+// RUN: %clang -emit-llvm -S -o - %s | FileCheck %s
+
+// CHECK: foo
+__kernel void foo(__global float4 *f) {
   *f = cos(*f);
 }

libclc/test/cross.cl

@@ -1,3 +1,6 @@
-__kernel void foo(float4 *f) {
+// RUN: %clang -emit-llvm -S -o - %s | FileCheck %s
+
+// CHECK: foo
+__kernel void foo(__global float4 *f) {
   *f = cross(f[0], f[1]);
 }

libclc/test/fabs.cl

@@ -1,3 +1,6 @@
-__kernel void foo(float *f) {
+// RUN: %clang -emit-llvm -S -o - %s | FileCheck %s
+
+// CHECK: foo
+__kernel void foo(__global float *f) {
   *f = fabs(*f);
 }

libclc/test/get_group_id.cl

@@ -1,3 +1,6 @@
-__kernel void foo(int *i) {
+// RUN: %clang -emit-llvm -S -o - %s | FileCheck %s
+
+// CHECK: foo
+__kernel void foo(__global int *i) {
   i[get_group_id(0)] = 1;
 }

libclc/test/lit.cfg.py

@@ -0,0 +1,57 @@
+import os
+
+import lit.formats
+import lit.util
+
+from lit.llvm import llvm_config
+from lit.llvm.subst import ToolSubst
+from lit.llvm.subst import FindTool
+
+# Configuration file for the 'lit' test runner.
+
+# name: The name of this test suite.
+config.name = "libclc"
+
+# testFormat: The test format to use to interpret tests.
+config.test_format = lit.formats.ShTest(not llvm_config.use_lit_shell)
+
+# suffixes: A list of file extensions to treat as test files.
+config.suffixes = [
+    ".cl",
+]
+
+# test_source_root: The root path where tests are located.
+config.test_source_root = os.path.join(os.path.dirname(__file__))
+
+# test_exec_root: The root path where tests should be run.
+config.test_exec_root = os.path.join(config.test_run_dir, "test")
+
+llvm_config.use_default_substitutions()
+
+target = lit_config.params.get("target", "")
+builtins = lit_config.params.get("builtins", "")
+
+clang_flags = [
+    "-fno-builtin",
+    "-target",
+    target,
+    "-Xclang",
+    "-mlink-builtin-bitcode",
+    "-Xclang",
+    os.path.join(config.libclc_lib_dir, builtins),
+    "-nogpulib",
+]
+
+cpu = lit_config.params.get("cpu", "")
+if cpu:
+    clang_flags.append(f"-mcpu={cpu}")
+
+llvm_config.use_clang(additional_flags=clang_flags)
+
+tools = [
+    "llvm-dis",
+    "not",
+]
+tool_dirs = [config.llvm_tools_dir]
+
+llvm_config.add_tool_substitutions(tools, tool_dirs)

libclc/test/lit.site.cfg.py.in

@@ -0,0 +1,24 @@
+@LIT_SITE_CFG_IN_HEADER@
+
+import sys
+
+config.llvm_src_root = path(r"@LLVM_SOURCE_DIR@")
+config.llvm_obj_root = path(r"@LLVM_BINARY_DIR@")
+config.llvm_tools_dir = lit_config.substitute(path(r"@LLVM_TOOLS_DIR@"))
+config.llvm_libs_dir = lit_config.substitute(path(r"@LLVM_LIBS_DIR@"))
+config.llvm_shlib_dir = lit_config.substitute(path(r"@SHLIBDIR@"))
+config.lit_tools_dir = path(r"@LLVM_LIT_TOOLS_DIR@")
+config.host_triple = "@LLVM_HOST_TRIPLE@"
+config.target_triple = "@LLVM_TARGET_TRIPLE@"
+config.host_arch = "@HOST_ARCH@"
+config.python_executable = "@Python3_EXECUTABLE@"
+config.libclc_src_dir = path(r"@LIBCLC_SOURCE_DIR@")
+config.libclc_lib_dir = path(r"@LIBCLC_BINARY_DIR@")
+config.test_run_dir = path(r"@LIBCLC_BINARY_DIR@")
+
+import lit.llvm
+lit.llvm.initialize(lit_config, config)
+
+# Let the main config do the real work.
+lit_config.load_config(
+    config, os.path.join(config.libclc_src_dir, "test/lit.cfg.py"))

libclc/test/rsqrt.cl

@@ -1,6 +1,10 @@
-#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+// RUN: %clang -emit-llvm -S %s
 
-__kernel void foo(float4 *x, double4 *y) {
+#if defined(cl_khr_fp64)
+
+__kernel void foo(__global float4 *x, __global double4 *y) {
   x[1] = rsqrt(x[0]);
   y[1] = rsqrt(y[0]);
 }
+
+#endif

libclc/test/subsat.cl

@@ -1,19 +1,22 @@
-__kernel void test_subsat_char(char *a, char x, char y) {
+// RUN: %clang -emit-llvm -S -o - %s | FileCheck %s
+
+// CHECK: test_subsat_char
+__kernel void test_subsat_char(__global char *a, char x, char y) {
   *a = sub_sat(x, y);
   return;
 }
 
-__kernel void test_subsat_uchar(uchar *a, uchar x, uchar y) {
+__kernel void test_subsat_uchar(__global uchar *a, uchar x, uchar y) {
   *a = sub_sat(x, y);
   return;
 }
 
-__kernel void test_subsat_long(long *a, long x, long y) {
+__kernel void test_subsat_long(__global long *a, long x, long y) {
   *a = sub_sat(x, y);
   return;
 }
 
-__kernel void test_subsat_ulong(ulong *a, ulong x, ulong y) {
+__kernel void test_subsat_ulong(__global ulong *a, ulong x, ulong y) {
   *a = sub_sat(x, y);
   return;
 }