[libclc] Add initial LIT tests

.ci/compute_projects.py

@@ -62,7 +62,6 @@
         "mlir",
         "polly",
         "flang",
-        "libclc",
         "openmp",
     },
 }
@@ -75,7 +74,7 @@
 # This mapping describes runtimes that should be tested when the key project is
 # touched.
 DEPENDENT_RUNTIMES_TO_TEST = {
-    "clang": {"compiler-rt"},
+    "clang": {"compiler-rt", "libclc"},
     "clang-tools-extra": {"libc"},
     "libc": {"libc"},
     ".ci": {"compiler-rt", "libc"},
@@ -113,6 +112,7 @@
     "cross-project-tests",
     "flang",
     "libc",
+    "libclc",
     "lldb",
     "openmp",
     "polly",
@@ -132,14 +132,15 @@
     "lld": "check-lld",
     "flang": "check-flang",
     "libc": "check-libc",
+    "libclc": "check-libclc",
     "lld": "check-lld",
     "lldb": "check-lldb",
     "mlir": "check-mlir",
     "openmp": "check-openmp",
     "polly": "check-polly",
 }
 
-RUNTIMES = {"libcxx", "libcxxabi", "libunwind", "compiler-rt", "libc"}
+RUNTIMES = {"libcxx", "libcxxabi", "libunwind", "compiler-rt", "libc", "libclc"}
 
 
 def _add_dependencies(projects: Set[str], runtimes: Set[str]) -> Set[str]:

clang/include/clang/Basic/DiagnosticDriverKinds.td

@@ -392,6 +392,8 @@ def warn_drv_fraw_string_literals_in_cxx11 : Warning<
   "ignoring '-f%select{no-|}0raw-string-literals', which is only valid for C and C++ standards before C++11">,
   InGroup<UnusedCommandLineArgument>;
 
+def err_drv_libclc_not_found : Error<"no libclc library '%0' found in the clang resource directory">;
+
 def err_drv_invalid_malign_branch_EQ : Error<
   "invalid argument '%0' to -malign-branch=; each element must be one of: %1">;
 

clang/include/clang/Driver/CommonArgs.h

@@ -215,6 +215,9 @@ void addOpenMPDeviceRTL(const Driver &D, const llvm::opt::ArgList &DriverArgs,
                         StringRef BitcodeSuffix, const llvm::Triple &Triple,
                         const ToolChain &HostTC);
 
+void addOpenCLBuiltinsLib(const Driver &D, const llvm::opt::ArgList &DriverArgs,
+                          llvm::opt::ArgStringList &CC1Args);
+
 void addOutlineAtomicsArgs(const Driver &D, const ToolChain &TC,
                            const llvm::opt::ArgList &Args,
                            llvm::opt::ArgStringList &CmdArgs,

clang/include/clang/Driver/Options.td

@@ -1425,6 +1425,8 @@ def openacc_macro_override_EQ
 
 // End Clang specific/exclusive options for OpenACC.
 
+def libclc_lib_EQ : Joined<["--"], "libclc-lib=">, Group<opencl_Group>,
+  HelpText<"Namespec of libclc OpenCL bitcode library to link">;
 def libomptarget_amdgpu_bc_path_EQ : Joined<["--"], "libomptarget-amdgpu-bc-path=">, Group<i_Group>,
   HelpText<"Path to libomptarget-amdgcn bitcode library">;
 def libomptarget_amdgcn_bc_path_EQ : Joined<["--"], "libomptarget-amdgcn-bc-path=">, Group<i_Group>,

clang/lib/Driver/ToolChains/AMDGPU.cpp

@@ -795,6 +795,8 @@ void AMDGPUToolChain::addClangTargetOptions(
     CC1Args.push_back("-fvisibility=hidden");
     CC1Args.push_back("-fapply-global-visibility-to-externs");
   }
+
+  addOpenCLBuiltinsLib(getDriver(), DriverArgs, CC1Args);
 }
 
 void AMDGPUToolChain::addClangWarningOptions(ArgStringList &CC1Args) const {

clang/lib/Driver/ToolChains/CommonArgs.cpp

@@ -2942,6 +2942,43 @@ void tools::addHIPRuntimeLibArgs(const ToolChain &TC, Compilation &C,
   }
 }
 
+void tools::addOpenCLBuiltinsLib(const Driver &D,
+                                 const llvm::opt::ArgList &DriverArgs,
+                                 llvm::opt::ArgStringList &CC1Args) {
+  // Check whether user specifies a libclc bytecode library
+  const Arg *A = DriverArgs.getLastArg(options::OPT_libclc_lib_EQ);
+  if (!A)
+    return;
+
+  // Find device libraries in <LLVM_DIR>/lib/clang/<ver>/lib/libclc/
+  SmallString<128> LibclcPath(D.ResourceDir);
+  llvm::sys::path::append(LibclcPath, "lib", "libclc");
+
+  // If the namespec is of the form :filename, search for that file.
+  StringRef LibclcNamespec(A->getValue());
+  bool FilenameSearch = LibclcNamespec.consume_front(":");
+  SmallString<128> LibclcTargetFile(LibclcNamespec);
+
+  if (FilenameSearch && llvm::sys::fs::exists(LibclcTargetFile)) {
+    CC1Args.push_back("-mlink-builtin-bitcode");
+    CC1Args.push_back(DriverArgs.MakeArgString(LibclcTargetFile));
+  } else {
+    // Search the library paths for the file
+    if (!FilenameSearch)
+      LibclcTargetFile += ".bc";
+
+    llvm::sys::path::append(LibclcPath, LibclcTargetFile);
+    if (llvm::sys::fs::exists(LibclcPath)) {
+      CC1Args.push_back("-mlink-builtin-bitcode");
+      CC1Args.push_back(DriverArgs.MakeArgString(LibclcPath));
+    } else {
+      // Since the user requested a library, if we haven't one then report an
+      // error.
+      D.Diag(diag::err_drv_libclc_not_found) << LibclcTargetFile;
+    }
+  }
+}
+
 void tools::addOutlineAtomicsArgs(const Driver &D, const ToolChain &TC,
                                   const llvm::opt::ArgList &Args,
                                   llvm::opt::ArgStringList &CmdArgs,

clang/test/Driver/opencl-libclc.cl

@@ -0,0 +1,9 @@
+// RUN: %clang -### -target amdgcn-amd-amdhsa --no-offloadlib --libclc-lib=:%S/Inputs/libclc/libclc.bc %s 2>&1 | FileCheck %s
+// RUN: %clang -### -target amdgcn-amd-amdhsa --no-offloadlib --libclc-lib=:%S/Inputs/libclc/subdir/libclc.bc %s 2>&1 | FileCheck %s --check-prefix CHECK-SUBDIR
+
+// RUN: not %clang -### -target amdgcn-amd-amdhsa --no-offloadlib --libclc-lib=:%S/Inputs/libclc/subdir/not-here.bc %s 2>&1 | FileCheck %s --check-prefix CHECK-ERROR
+
+// CHECK: -mlink-builtin-bitcode{{.*}}Inputs{{/|\\\\}}libclc{{/|\\\\}}libclc.bc
+// CHECK-SUBDIR: -mlink-builtin-bitcode{{.*}}Inputs{{/|\\\\}}libclc{{/|\\\\}}subdir{{/|\\\\}}libclc.bc
+
+// CHECK-ERROR: no libclc library{{.*}}not-here.bc' found in the clang resource directory

libclc/CMakeLists.txt

@@ -63,6 +63,9 @@ if( LIBCLC_STANDALONE_BUILD OR CMAKE_SOURCE_DIR STREQUAL CMAKE_CURRENT_SOURCE_DI
       set( ${tool}_target )
     endforeach()
   endif()
+
+  # Setup the paths where libclc runtimes should be stored.
+  set( LIBCLC_OUTPUT_LIBRARY_DIR ${CMAKE_CURRENT_BINARY_DIR} )
 else()
   # In-tree configuration
   set( LIBCLC_STANDALONE_BUILD FALSE )
@@ -82,10 +85,14 @@ else()
     get_host_tool_path( llvm-link LLVM_LINK llvm-link_exe llvm-link_target )
     get_host_tool_path( opt OPT opt_exe opt_target )
   endif()
-endif()
 
-# Setup the paths where libclc runtimes should be stored.
-set( LIBCLC_OUTPUT_LIBRARY_DIR ${CMAKE_CURRENT_BINARY_DIR} )
+  # Setup the paths where libclc runtimes should be stored. By default, in an
+  # in-tree build we place the libraries in clang's resource driectory.
+  get_clang_resource_dir( LIBCLC_OUTPUT_DIR PREFIX ${LLVM_LIBRARY_OUTPUT_INTDIR}/.. )
+
+  # Note we do not adhere to LLVM_ENABLE_PER_TARGET_RUNTIME_DIR.
+  set( LIBCLC_OUTPUT_LIBRARY_DIR ${LIBCLC_OUTPUT_DIR}/lib/libclc )
+endif()
 
 if( EXISTS ${LIBCLC_CUSTOM_LLVM_TOOLS_BINARY_DIR} )
   message( WARNING "Using custom LLVM tools to build libclc: "
@@ -487,3 +494,7 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} )
     )
   endforeach( d )
 endforeach( t )
+
+if( NOT LIBCLC_STANDALONE_BUILD )
+  add_subdirectory( test )
+endif()

libclc/test/.clang-format

@@ -0,0 +1 @@
+DisableFormat: true

libclc/test/CMakeLists.txt

@@ -0,0 +1,35 @@
+set( LIBCLC_TEST_SOURCE_DIR ${CMAKE_CURRENT_SOURCE_DIR} )
+
+set( LIBCLC_TEST_TARGETS_ALL
+  amdgcn-mesa-mesa3d
+)
+
+foreach( target IN LISTS LIBCLC_TEST_TARGETS_ALL )
+  # If we haven't built this libclc target, don't build the tests
+  if( NOT TARGET prepare-${target} )
+    message( WARNING "libclc tests require target ${target}. Tests will not be built" )
+    # Add a dummy target
+    add_custom_target( check-libclc )
+    return()
+  endif()
+
+  list( APPEND LIBCLC_TEST_DEPS prepare-${target} )
+endforeach()
+
+list( APPEND LIBCLC_TEST_DEPS
+  ${clang_target}
+  FileCheck
+)
+
+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
+)
+
+add_lit_testsuite( check-libclc
+  "Running libclc regression tests"
+  ${CMAKE_CURRENT_BINARY_DIR}
+  DEPENDS ${LIBCLC_TEST_DEPS}
+)

libclc/test/geometric/cross.cl

@@ -0,0 +1,51 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+// RUN: %clang -target amdgcn-mesa-mesa3d -O1 -emit-llvm -S -o - --libclc-lib=tahiti-amdgcn-mesa-mesa3d -fno-builtin --no-offloadlib %s | FileCheck %s
+
+// CHECK-LABEL: define protected amdgpu_kernel void @foo(
+// CHECK-SAME: ptr addrspace(1) noundef align 16 captures(none) [[F:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] !kernel_arg_addr_space [[META6:![0-9]+]] !kernel_arg_access_qual [[META7:![0-9]+]] !kernel_arg_type [[META8:![0-9]+]] !kernel_arg_base_type [[META9:![0-9]+]] !kernel_arg_type_qual [[META10:![0-9]+]] {
+// CHECK-NEXT:  [[ENTRY:.*:]]
+// CHECK-NEXT:    [[TMP0:%.*]] = load <4 x float>, ptr addrspace(1) [[F]], align 16, !tbaa [[TBAA11:![0-9]+]]
+// CHECK-NEXT:    [[ARRAYIDX1_I:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(1) [[F]], i64 16
+// CHECK-NEXT:    [[TMP1:%.*]] = load <4 x float>, ptr addrspace(1) [[ARRAYIDX1_I]], align 16, !tbaa [[TBAA11]]
+// CHECK-NEXT:    [[TMP2:%.*]] = extractelement <4 x float> [[TMP0]], i64 1
+// CHECK-NEXT:    [[TMP3:%.*]] = extractelement <4 x float> [[TMP1]], i64 2
+// CHECK-NEXT:    [[TMP4:%.*]] = extractelement <4 x float> [[TMP0]], i64 2
+// CHECK-NEXT:    [[TMP5:%.*]] = extractelement <4 x float> [[TMP1]], i64 1
+// CHECK-NEXT:    [[TMP6:%.*]] = fneg float [[TMP5]]
+// CHECK-NEXT:    [[NEG_I_I:%.*]] = fmul float [[TMP4]], [[TMP6]]
+// CHECK-NEXT:    [[TMP7:%.*]] = tail call float @llvm.fmuladd.f32(float [[TMP2]], float [[TMP3]], float [[NEG_I_I]])
+// CHECK-NEXT:    [[TMP8:%.*]] = extractelement <4 x float> [[TMP1]], i64 0
+// CHECK-NEXT:    [[TMP9:%.*]] = extractelement <4 x float> [[TMP0]], i64 0
+// CHECK-NEXT:    [[TMP10:%.*]] = fneg float [[TMP3]]
+// CHECK-NEXT:    [[NEG3_I_I:%.*]] = fmul float [[TMP9]], [[TMP10]]
+// CHECK-NEXT:    [[TMP11:%.*]] = tail call float @llvm.fmuladd.f32(float [[TMP4]], float [[TMP8]], float [[NEG3_I_I]])
+// CHECK-NEXT:    [[TMP12:%.*]] = fneg float [[TMP8]]
+// CHECK-NEXT:    [[NEG6_I_I:%.*]] = fmul float [[TMP2]], [[TMP12]]
+// CHECK-NEXT:    [[TMP13:%.*]] = tail call float @llvm.fmuladd.f32(float [[TMP9]], float [[TMP5]], float [[NEG6_I_I]])
+// CHECK-NEXT:    [[TMP14:%.*]] = insertelement <4 x float> <float poison, float poison, float poison, float 0.000000e+00>, float [[TMP7]], i64 0
+// CHECK-NEXT:    [[TMP15:%.*]] = insertelement <4 x float> [[TMP14]], float [[TMP11]], i64 1
+// CHECK-NEXT:    [[VECINIT8_I_I:%.*]] = insertelement <4 x float> [[TMP15]], float [[TMP13]], i64 2
+// CHECK-NEXT:    store <4 x float> [[VECINIT8_I_I]], ptr addrspace(1) [[F]], align 16, !tbaa [[TBAA11]]
+// CHECK-NEXT:    ret void
+//
+__kernel void foo(__global float4 *f) {
+  *f = cross(f[0], f[1]);
+}
+//.
+// CHECK: [[META6]] = !{i32 1}
+// CHECK: [[META7]] = !{!"none"}
+// CHECK: [[META8]] = !{!"float4*"}
+// CHECK: [[META9]] = !{!"float __attribute__((ext_vector_type(4)))*"}
+// CHECK: [[META10]] = !{!""}
+// CHECK: [[TBAA11]] = !{[[META12:![0-9]+]], [[META12]], i64 0}
+// CHECK: [[META12]] = !{!"omnipotent char", [[META13:![0-9]+]], i64 0}
+// CHECK: [[META13]] = !{!"Simple C/C++ TBAA"}
+//.

libclc/test/integer/add_sat.cl

@@ -0,0 +1,32 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+// RUN: %clang -target amdgcn-mesa-mesa3d -O1 -emit-llvm -S -o - --libclc-lib=tahiti-amdgcn-mesa-mesa3d -fno-builtin --no-offloadlib %s | FileCheck %s
+
+// CHECK-LABEL: define protected amdgpu_kernel void @foo(
+// CHECK-SAME: ptr addrspace(1) noundef writeonly align 1 captures(none) initializes((0, 1)) [[A:%.*]], ptr addrspace(1) noundef readonly align 1 captures(none) [[B:%.*]], ptr addrspace(1) noundef readonly align 1 captures(none) [[C:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] !kernel_arg_addr_space [[META6:![0-9]+]] !kernel_arg_access_qual [[META7:![0-9]+]] !kernel_arg_type [[META8:![0-9]+]] !kernel_arg_base_type [[META8]] !kernel_arg_type_qual [[META9:![0-9]+]] {
+// CHECK-NEXT:  [[ENTRY:.*:]]
+// CHECK-NEXT:    [[TMP0:%.*]] = load i8, ptr addrspace(1) [[B]], align 1, !tbaa [[TBAA10:![0-9]+]]
+// CHECK-NEXT:    [[TMP1:%.*]] = load i8, ptr addrspace(1) [[C]], align 1, !tbaa [[TBAA10]]
+// CHECK-NEXT:    [[ELT_SAT_I_I:%.*]] = tail call noundef i8 @llvm.sadd.sat.i8(i8 [[TMP0]], i8 [[TMP1]])
+// CHECK-NEXT:    store i8 [[ELT_SAT_I_I]], ptr addrspace(1) [[A]], align 1, !tbaa [[TBAA10]]
+// CHECK-NEXT:    ret void
+//
+__kernel void foo(__global char *a, __global char *b, __global char *c) {
+  *a = add_sat(*b, *c);
+}
+//.
+// CHECK: [[META6]] = !{i32 1, i32 1, i32 1}
+// CHECK: [[META7]] = !{!"none", !"none", !"none"}
+// CHECK: [[META8]] = !{!"char*", !"char*", !"char*"}
+// CHECK: [[META9]] = !{!"", !"", !""}
+// CHECK: [[TBAA10]] = !{[[META11:![0-9]+]], [[META11]], i64 0}
+// CHECK: [[META11]] = !{!"omnipotent char", [[META12:![0-9]+]], i64 0}
+// CHECK: [[META12]] = !{!"Simple C/C++ TBAA"}
+//.