[libclc] [llvm] [WIP] Libclc tests (PR #87989)
Fraser Cormack via cfe-commits
cfe-commits at lists.llvm.org
Mon Apr 8 06:26:57 PDT 2024
https://github.com/frasercrmck created https://github.com/llvm/llvm-project/pull/87989
None
>From b87c2862491aed03f883d6b124a845b0d7c47699 Mon Sep 17 00:00:00 2001
From: Fraser Cormack <fraser at codeplay.com>
Date: Wed, 3 Apr 2024 16:09:30 +0100
Subject: [PATCH 1/2] [libclc] Refactor build system to allow in-tree builds
The previous build system was adding custom "OpenCL" and "LLVM IR"
languages in CMake to build the builtin libraries. This was making it
harder to build in-tree because the tool binaries needed to be present
at configure time.
This commit refactors the build system to use custom commands to build
the bytecode files one by one, and link them all together into the final
bytecode library. It also enables in-tree builds by aliasing the
clang/llvm-link/etc. tool targets to internal targets, which are
imported from the LLVM installation directory when building out of tree.
Diffing (with llvm-diff) all of the final bytecode libraries in an
out-of-tree configuration against those built using the current tip
system shows no changes. Note that there are textual changes to metadata
IDs which confuse regular diff, and that llvm-diff 14 and below may show
false-positives.
This commit also removes a file listed in one of the SOURCEs which
didn't exist and which was preventing the use of
ENABLE_RUNTIME_SUBNORMAL when configuring CMake.
---
libclc/CMakeLists.txt | 233 ++++++++++--------
libclc/cmake/CMakeCLCCompiler.cmake.in | 9 -
libclc/cmake/CMakeCLCInformation.cmake | 12 -
libclc/cmake/CMakeDetermineCLCCompiler.cmake | 18 --
.../cmake/CMakeDetermineLLAsmCompiler.cmake | 24 --
libclc/cmake/CMakeLLAsmCompiler.cmake.in | 10 -
libclc/cmake/CMakeLLAsmInformation.cmake | 12 -
libclc/cmake/CMakeTestCLCCompiler.cmake | 56 -----
libclc/cmake/CMakeTestLLAsmCompiler.cmake | 56 -----
libclc/cmake/modules/AddLibclc.cmake | 153 ++++++++++++
libclc/generic/lib/SOURCES | 1 -
llvm/tools/CMakeLists.txt | 3 +
12 files changed, 293 insertions(+), 294 deletions(-)
delete mode 100644 libclc/cmake/CMakeCLCCompiler.cmake.in
delete mode 100644 libclc/cmake/CMakeCLCInformation.cmake
delete mode 100644 libclc/cmake/CMakeDetermineCLCCompiler.cmake
delete mode 100644 libclc/cmake/CMakeDetermineLLAsmCompiler.cmake
delete mode 100644 libclc/cmake/CMakeLLAsmCompiler.cmake.in
delete mode 100644 libclc/cmake/CMakeLLAsmInformation.cmake
delete mode 100644 libclc/cmake/CMakeTestCLCCompiler.cmake
delete mode 100644 libclc/cmake/CMakeTestLLAsmCompiler.cmake
create mode 100644 libclc/cmake/modules/AddLibclc.cmake
diff --git a/libclc/CMakeLists.txt b/libclc/CMakeLists.txt
index 8750a65a717f9b..e509a61f51b927 100644
--- a/libclc/CMakeLists.txt
+++ b/libclc/CMakeLists.txt
@@ -4,6 +4,15 @@ project( libclc VERSION 0.2.0 LANGUAGES CXX C)
set(CMAKE_CXX_STANDARD 17)
+# Add path for custom modules
+list( INSERT CMAKE_MODULE_PATH 0 "${PROJECT_SOURCE_DIR}/cmake/modules" )
+
+set( LIBCLC_SOURCE_DIR ${CMAKE_CURRENT_SOURCE_DIR} )
+set( LIBCLC_BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR} )
+set( LIBCLC_OBJFILE_DIR ${LIBCLC_BINARY_DIR}/obj.libclc.dir )
+
+include( AddLibclc )
+
include( GNUInstallDirs )
set_property(DIRECTORY APPEND PROPERTY CMAKE_CONFIGURE_DEPENDS
amdgcn-amdhsa/lib/SOURCES;
@@ -27,31 +36,51 @@ set( LIBCLC_TARGETS_TO_BUILD "all"
option( ENABLE_RUNTIME_SUBNORMAL "Enable runtime linking of subnormal support." OFF )
-find_package(LLVM REQUIRED HINTS "${LLVM_CMAKE_DIR}")
-include(AddLLVM)
+if( LIBCLC_STANDALONE_BUILD OR CMAKE_SOURCE_DIR STREQUAL CMAKE_CURRENT_SOURCE_DIR )
+ # Out-of-tree configuration
+ set( LIBCLC_STANDALONE_BUILD TRUE )
-message( STATUS "libclc LLVM version: ${LLVM_PACKAGE_VERSION}" )
+ find_package(LLVM REQUIRED HINTS "${LLVM_CMAKE_DIR}")
+ include(AddLLVM)
-if( LLVM_PACKAGE_VERSION VERSION_LESS LIBCLC_MIN_LLVM )
- message( FATAL_ERROR "libclc needs at least LLVM ${LIBCLC_MIN_LLVM}" )
-endif()
+ message( STATUS "libclc LLVM version: ${LLVM_PACKAGE_VERSION}" )
-find_program( LLVM_CLANG clang PATHS ${LLVM_TOOLS_BINARY_DIR} NO_DEFAULT_PATH )
-find_program( LLVM_AS llvm-as PATHS ${LLVM_TOOLS_BINARY_DIR} NO_DEFAULT_PATH )
-find_program( LLVM_LINK llvm-link PATHS ${LLVM_TOOLS_BINARY_DIR} NO_DEFAULT_PATH )
-find_program( LLVM_OPT opt PATHS ${LLVM_TOOLS_BINARY_DIR} NO_DEFAULT_PATH )
-find_program( LLVM_SPIRV llvm-spirv PATHS ${LLVM_TOOLS_BINARY_DIR} NO_DEFAULT_PATH )
+ if( LLVM_PACKAGE_VERSION VERSION_LESS LIBCLC_MIN_LLVM )
+ message( FATAL_ERROR "libclc needs at least LLVM ${LIBCLC_MIN_LLVM}" )
+ endif()
+
+ # Import required tools as targets
+ foreach( tool clang llvm-as llvm-link opt )
+ find_program( LLVM_TOOL_${tool} ${tool} PATHS ${LLVM_TOOLS_BINARY_DIR} NO_DEFAULT_PATH )
+ add_executable( libclc::${tool} IMPORTED GLOBAL )
+ set_target_properties( libclc::${tool} PROPERTIES IMPORTED_LOCATION ${LLVM_TOOL_${tool}} )
+ endforeach()
+else()
+ # In-tree configuration
+ set( LIBCLC_STANDALONE_BUILD FALSE )
+
+ set( LLVM_PACKAGE_VERSION ${LLVM_VERSION} )
-# Print toolchain
-message( STATUS "libclc toolchain - clang: ${LLVM_CLANG}" )
-message( STATUS "libclc toolchain - llvm-as: ${LLVM_AS}" )
-message( STATUS "libclc toolchain - llvm-link: ${LLVM_LINK}" )
-message( STATUS "libclc toolchain - opt: ${LLVM_OPT}" )
-message( STATUS "libclc toolchain - llvm-spirv: ${LLVM_SPIRV}" )
-if( NOT LLVM_CLANG OR NOT LLVM_OPT OR NOT LLVM_AS OR NOT LLVM_LINK )
+ # Note that we check this later (for both build types) but we can provide a
+ # more useful error message when built in-tree. We assume that LLVM tools are
+ # always available so don't warn here.
+ if( NOT clang IN_LIST LLVM_ENABLE_PROJECTS )
+ message(FATAL_ERROR "Clang is not enabled, but is required to build libclc in-tree")
+ endif()
+
+ foreach( tool clang llvm-as llvm-link opt )
+ add_executable(libclc::${tool} ALIAS ${tool})
+ endforeach()
+endif()
+
+if( NOT TARGET libclc::clang OR NOT TARGET libclc::opt
+ OR NOT TARGET libclc::llvm-as OR NOT TARGET libclc::llvm-link )
message( FATAL_ERROR "libclc toolchain incomplete!" )
endif()
+# llvm-spirv is an optional dependency, used to build spirv-* targets.
+find_program( LLVM_SPIRV llvm-spirv PATHS ${LLVM_TOOLS_BINARY_DIR} NO_DEFAULT_PATH )
+
# List of all targets. Note that some are added dynamically below.
set( LIBCLC_TARGETS_ALL
amdgcn--
@@ -90,24 +119,9 @@ if( "spirv-mesa3d-" IN_LIST LIBCLC_TARGETS_TO_BUILD OR "spirv64-mesa3d-" IN_LIST
endif()
endif()
-set( CMAKE_MODULE_PATH ${PROJECT_SOURCE_DIR}/cmake )
-set( CMAKE_CLC_COMPILER ${LLVM_CLANG} )
-set( CMAKE_CLC_ARCHIVE ${LLVM_LINK} )
-set( CMAKE_LLAsm_PREPROCESSOR ${LLVM_CLANG} )
-set( CMAKE_LLAsm_COMPILER ${LLVM_AS} )
-set( CMAKE_LLAsm_ARCHIVE ${LLVM_LINK} )
-
# Construct LLVM version define
set( LLVM_VERSION_DEFINE "-DHAVE_LLVM=0x${LLVM_VERSION_MAJOR}0${LLVM_VERSION_MINOR}" )
-
-# LLVM 13 enables standard includes by default
-if( LLVM_PACKAGE_VERSION VERSION_GREATER_EQUAL 13.0.0 )
- set( CMAKE_LLAsm_FLAGS "${CMAKE_LLAsm_FLAGS} -cl-no-stdinc" )
- set( CMAKE_CLC_FLAGS "${CMAKE_CLC_FLAGS} -cl-no-stdinc" )
-endif()
-
-enable_language( CLC LLAsm )
# This needs to be set before any target that needs it
# We need to use LLVM_INCLUDE_DIRS here, because if we are linking to an
# llvm build directory, this includes $src/llvm/include which is where all the
@@ -122,7 +136,7 @@ set(LLVM_LINK_COMPONENTS
IRReader
Support
)
-if(CMAKE_SOURCE_DIR STREQUAL CMAKE_CURRENT_SOURCE_DIR)
+if( LIBCLC_STANDALONE_BUILD )
add_llvm_executable( prepare_builtins utils/prepare-builtins.cpp )
else()
add_llvm_utility( prepare_builtins utils/prepare-builtins.cpp )
@@ -167,12 +181,15 @@ install( FILES ${CMAKE_CURRENT_BINARY_DIR}/libclc.pc DESTINATION "${CMAKE_INSTAL
install( DIRECTORY generic/include/clc DESTINATION "${CMAKE_INSTALL_INCLUDEDIR}" )
if( ENABLE_RUNTIME_SUBNORMAL )
- add_library( subnormal_use_default STATIC
- generic/lib/subnormal_use_default.ll )
- add_library( subnormal_disable STATIC
- generic/lib/subnormal_disable.ll )
- install( TARGETS subnormal_use_default subnormal_disable ARCHIVE
- DESTINATION "${CMAKE_INSTALL_DATADIR}/clc" )
+ foreach( file subnormal_use_default subnormal_disable )
+ link_bc(
+ OUTPUT ${LIBCLC_OBJFILE_DIR}/${file}.bc
+ INPUTS ${PROJECT_SOURCE_DIR}/generic/lib/${file}.ll
+ )
+ add_custom_target( ${file}.bc ALL DEPENDS ${LIBCLC_OBJFILE_DIR}/${file}.bc )
+ install( FILES ${LIBCLC_OBJFILE_DIR}/${file}.bc ARCHIVE
+ DESTINATION "${CMAKE_INSTALL_DATADIR}/clc" )
+ endforeach()
endif()
find_package( Python3 REQUIRED COMPONENTS Interpreter )
@@ -232,19 +249,19 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} )
# Add the generated convert.cl here to prevent adding the one listed in
# SOURCES
+ set( objects ) # A "set" of already-added input files
+ set( rel_files ) # Source directory input files, relative to the root dir
+ set( gen_files ) # Generated binary input files, relative to the binary dir
if( NOT ${ARCH} STREQUAL "spirv" AND NOT ${ARCH} STREQUAL "spirv64" )
if( NOT ENABLE_RUNTIME_SUBNORMAL AND NOT ${ARCH} STREQUAL "clspv" AND
NOT ${ARCH} STREQUAL "clspv64" )
- set( rel_files convert.cl )
- set( objects convert.cl )
+ list( APPEND gen_files convert.cl )
+ list( APPEND objects convert.cl )
list( APPEND rel_files generic/lib/subnormal_use_default.ll )
elseif(${ARCH} STREQUAL "clspv" OR ${ARCH} STREQUAL "clspv64")
- set( rel_files clspv-convert.cl )
- set( objects clspv-convert.cl )
+ list( APPEND gen_files clspv-convert.cl )
+ list( APPEND objects clspv-convert.cl )
endif()
- else()
- set( rel_files )
- set( objects )
endif()
foreach( l ${source_list} )
@@ -252,46 +269,35 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} )
string( REPLACE "\n" ";" file_list ${file_list} )
get_filename_component( dir ${l} DIRECTORY )
foreach( f ${file_list} )
- list( FIND objects ${f} found )
- if( found EQUAL -1 )
+ # Only add each file once, so that targets can 'specialize' builtins
+ if( NOT ${f} IN_LIST objects )
list( APPEND objects ${f} )
list( APPEND rel_files ${dir}/${f} )
- # FIXME: This should really go away
- file( TO_CMAKE_PATH ${PROJECT_SOURCE_DIR}/${dir}/${f} src_loc )
- get_filename_component( fdir ${src_loc} DIRECTORY )
-
- set_source_files_properties( ${dir}/${f}
- PROPERTIES COMPILE_FLAGS "-I ${fdir}" )
endif()
endforeach()
endforeach()
foreach( d ${${t}_devices} )
- # Some targets don't have a specific GPU to target
- if( ${d} STREQUAL "none" OR ${ARCH} STREQUAL "spirv" OR ${ARCH} STREQUAL "spirv64" )
- set( mcpu )
- set( arch_suffix "${t}" )
- else()
- set( mcpu "-mcpu=${d}" )
- set( arch_suffix "${d}-${t}" )
+ get_libclc_device_info(
+ TRIPLE ${t}
+ DEVICE ${d}
+ CPU cpu
+ ARCH_SUFFIX arch_suffix
+ CLANG_TRIPLE clang_triple
+ )
+
+ set( mcpu )
+ if( NOT "${cpu}" STREQUAL "" )
+ set( mcpu "-mcpu=${cpu}" )
endif()
+
message( STATUS " device: ${d} ( ${${d}_aliases} )" )
- if ( ${ARCH} STREQUAL "spirv" OR ${ARCH} STREQUAL "spirv64" )
- if( ${ARCH} STREQUAL "spirv" )
- set( t "spir--" )
- else()
- set( t "spir64--" )
- endif()
+ if ( ARCH STREQUAL spirv OR ARCH STREQUAL spirv64 )
set( build_flags -O0 -finline-hint-functions )
set( opt_flags )
set( spvflags --spirv-max-version=1.1 )
- elseif( ${ARCH} STREQUAL "clspv" )
- set( t "spir--" )
- set( build_flags "-Wno-unknown-assumption")
- set( opt_flags -O3 )
- elseif( ${ARCH} STREQUAL "clspv64" )
- set( t "spir64--" )
+ elseif( ARCH STREQUAL clspv OR ARCH STREQUAL clspv64 )
set( build_flags "-Wno-unknown-assumption")
set( opt_flags -O3 )
else()
@@ -299,32 +305,67 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} )
set( opt_flags -O3 )
endif()
- set( builtins_link_lib_tgt builtins.link.${arch_suffix} )
- add_library( ${builtins_link_lib_tgt} STATIC ${rel_files} )
- # Make sure we depend on the pseudo target to prevent
- # multiple invocations
- add_dependencies( ${builtins_link_lib_tgt} generate_convert.cl )
- add_dependencies( ${builtins_link_lib_tgt} clspv-generate_convert.cl )
- # CMake will turn this include into absolute path
- target_include_directories( ${builtins_link_lib_tgt} PRIVATE
- "generic/include" )
- target_compile_definitions( ${builtins_link_lib_tgt} PRIVATE
- "__CLC_INTERNAL" )
- string( TOUPPER "-DCLC_${ARCH}" CLC_TARGET_DEFINE )
- target_compile_definitions( ${builtins_link_lib_tgt} PRIVATE
- ${CLC_TARGET_DEFINE} )
- target_compile_options( ${builtins_link_lib_tgt} PRIVATE -target
- ${t} ${mcpu} -fno-builtin -nostdlib ${build_flags} )
- set_target_properties( ${builtins_link_lib_tgt} PROPERTIES
- LINKER_LANGUAGE CLC )
+ set( LIBCLC_ARCH_OBJFILE_DIR "${LIBCLC_OBJFILE_DIR}/${arch_suffix}" )
+ file( MAKE_DIRECTORY ${LIBCLC_ARCH_OBJFILE_DIR} )
+
+ string( TOUPPER "CLC_${ARCH}" CLC_TARGET_DEFINE )
+
+ list( APPEND build_flags
+ -D__CLC_INTERNAL
+ -D${CLC_TARGET_DEFINE}
+ -I${PROJECT_SOURCE_DIR}/generic/include
+ # FIXME: Fix libclc to not require disabling this noisy warnings
+ -Wno-bitwise-conditional-parentheses
+ )
+
+ set( bytecode_files "" )
+ foreach( file IN LISTS gen_files rel_files )
+ # We need to take each file and produce an absolute input file, as well
+ # as a unique architecture-specific output file. We deal with a mix of
+ # different input files, which makes this trickier.
+ if( ${file} IN_LIST gen_files )
+ # Generated files are given just as file names, which we must make
+ # absolute to the binary directory.
+ set( input_file ${CMAKE_CURRENT_BINARY_DIR}/${file} )
+ set( output_file "${LIBCLC_ARCH_OBJFILE_DIR}/${file}.o" )
+ else()
+ # Other files are originally relative to each SOURCE file, which are
+ # then make relative to the libclc root directory. We must normalize
+ # the path (e.g., ironing out any ".."), then make it relative to the
+ # root directory again, and use that relative path component for the
+ # binary path.
+ get_filename_component( abs_path ${file} ABSOLUTE BASE_DIR ${PROJECT_SOURCE_DIR} )
+ file( RELATIVE_PATH root_rel_path ${PROJECT_SOURCE_DIR} ${abs_path} )
+ set( input_file ${PROJECT_SOURCE_DIR}/${file} )
+ set( output_file "${LIBCLC_ARCH_OBJFILE_DIR}/${root_rel_path}.o" )
+ endif()
+
+ get_filename_component( file_dir ${file} DIRECTORY )
+
+ compile_to_bc(
+ TRIPLE ${clang_triple}
+ INPUT ${input_file}
+ OUTPUT ${output_file}
+ EXTRA_OPTS "${mcpu}" -fno-builtin -nostdlib
+ "${build_flags}" -I${PROJECT_SOURCE_DIR}/${file_dir}
+ )
+ list(APPEND bytecode_files ${output_file})
+ endforeach()
+
+ set( builtins_link_lib_tgt builtins.link.${arch_suffix}.bc )
+
+ link_bc(
+ OUTPUT ${builtins_link_lib_tgt}
+ INPUTS ${bytecode_files}
+ )
set( obj_suffix ${arch_suffix}.bc )
set( builtins_opt_lib_tgt builtins.opt.${obj_suffix} )
# Add opt target
add_custom_command( OUTPUT ${builtins_opt_lib_tgt}
- COMMAND ${LLVM_OPT} ${opt_flags} -o ${builtins_opt_lib_tgt}
- $<TARGET_FILE:${builtins_link_lib_tgt}>
+ COMMAND libclc::opt ${opt_flags} -o ${builtins_opt_lib_tgt}
+ ${builtins_link_lib_tgt}
DEPENDS ${builtins_link_lib_tgt} )
add_custom_target( "opt.${obj_suffix}" ALL
DEPENDS ${builtins_opt_lib_tgt} )
@@ -345,7 +386,7 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} )
add_custom_target( "prepare-${obj_suffix}" ALL DEPENDS "${obj_suffix}" )
# nvptx-- targets don't include workitem builtins
- if( NOT ${t} MATCHES ".*ptx.*--$" )
+ if( NOT ${clang_triple} MATCHES ".*ptx.*--$" )
add_test( NAME external-calls-${obj_suffix}
COMMAND ./check_external_calls.sh ${CMAKE_CURRENT_BINARY_DIR}/${obj_suffix} ${LLVM_TOOLS_BINARY_DIR}
WORKING_DIRECTORY ${PROJECT_SOURCE_DIR} )
@@ -353,7 +394,7 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} )
install( FILES ${CMAKE_CURRENT_BINARY_DIR}/${obj_suffix} DESTINATION "${CMAKE_INSTALL_DATADIR}/clc" )
foreach( a ${${d}_aliases} )
- set( alias_suffix "${a}-${t}.bc" )
+ 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}" )
diff --git a/libclc/cmake/CMakeCLCCompiler.cmake.in b/libclc/cmake/CMakeCLCCompiler.cmake.in
deleted file mode 100644
index 2730b83d9e7d09..00000000000000
--- a/libclc/cmake/CMakeCLCCompiler.cmake.in
+++ /dev/null
@@ -1,9 +0,0 @@
-set(CMAKE_CLC_COMPILER "@CMAKE_CLC_COMPILER@")
-set(CMAKE_CLC_COMPILER_LOADED 1)
-
-set(CMAKE_CLC_SOURCE_FILE_EXTENSIONS cl)
-set(CMAKE_CLC_OUTPUT_EXTENSION .bc)
-set(CMAKE_CLC_OUTPUT_EXTENSION_REPLACE 1)
-set(CMAKE_STATIC_LIBRARY_PREFIX_CLC "")
-set(CMAKE_STATIC_LIBRARY_SUFFIX_CLC ".bc")
-set(CMAKE_CLC_COMPILER_ENV_VAR "CLC_COMPILER")
diff --git a/libclc/cmake/CMakeCLCInformation.cmake b/libclc/cmake/CMakeCLCInformation.cmake
deleted file mode 100644
index 95327e44397222..00000000000000
--- a/libclc/cmake/CMakeCLCInformation.cmake
+++ /dev/null
@@ -1,12 +0,0 @@
-if(NOT CMAKE_CLC_COMPILE_OBJECT)
- set(CMAKE_CLC_COMPILE_OBJECT
- "<CMAKE_CLC_COMPILER> <DEFINES> <INCLUDES> <FLAGS> -o <OBJECT> -c <SOURCE> -emit-llvm")
-endif()
-
-if(NOT CMAKE_CLC_CREATE_STATIC_LIBRARY)
- set(CMAKE_CLC_CREATE_STATIC_LIBRARY
- "<CMAKE_CLC_ARCHIVE> -o <TARGET> <OBJECTS>")
-endif()
-
-set(CMAKE_INCLUDE_FLAG_CLC "-I")
-set(CMAKE_DEPFILE_FLAGS_CLC "-MD -MT <DEP_TARGET> -MF <DEP_FILE>")
diff --git a/libclc/cmake/CMakeDetermineCLCCompiler.cmake b/libclc/cmake/CMakeDetermineCLCCompiler.cmake
deleted file mode 100644
index 94d85d9e666adb..00000000000000
--- a/libclc/cmake/CMakeDetermineCLCCompiler.cmake
+++ /dev/null
@@ -1,18 +0,0 @@
-include(${CMAKE_ROOT}/Modules/CMakeDetermineCompiler.cmake)
-
-if(NOT CMAKE_CLC_COMPILER)
- find_program(CMAKE_CLC_COMPILER NAMES clang)
-endif()
-mark_as_advanced(CMAKE_CLC_COMPILER)
-
-if(NOT CMAKE_CLC_ARCHIVE)
- find_program(CMAKE_CLC_ARCHIVE NAMES llvm-link)
-endif()
-mark_as_advanced(CMAKE_CLC_ARCHIVE)
-
-set(CMAKE_CLC_COMPILER_ENV_VAR "CLC_COMPILER")
-set(CMAKE_CLC_ARCHIVE_ENV_VAR "CLC_LINKER")
-find_file(clc_comp_in CMakeCLCCompiler.cmake.in PATHS ${CMAKE_ROOT}/Modules ${CMAKE_MODULE_PATH})
-# configure all variables set in this file
-configure_file(${clc_comp_in} ${CMAKE_PLATFORM_INFO_DIR}/CMakeCLCCompiler.cmake @ONLY)
-mark_as_advanced(clc_comp_in)
diff --git a/libclc/cmake/CMakeDetermineLLAsmCompiler.cmake b/libclc/cmake/CMakeDetermineLLAsmCompiler.cmake
deleted file mode 100644
index 1c424c79cb453e..00000000000000
--- a/libclc/cmake/CMakeDetermineLLAsmCompiler.cmake
+++ /dev/null
@@ -1,24 +0,0 @@
-include(${CMAKE_ROOT}/Modules/CMakeDetermineCompiler.cmake)
-
-if(NOT CMAKE_LLAsm_PREPROCESSOR)
- find_program(CMAKE_LLAsm_PREPROCESSOR NAMES clang)
-endif()
-mark_as_advanced(CMAKE_LLAsm_PREPROCESSOR)
-
-if(NOT CMAKE_LLAsm_COMPILER)
- find_program(CMAKE_LLAsm_COMPILER NAMES llvm-as)
-endif()
-mark_as_advanced(CMAKE_LLAsm_ASSEMBLER)
-
-if(NOT CMAKE_LLAsm_ARCHIVE)
- find_program(CMAKE_LLAsm_ARCHIVE NAMES llvm-link)
-endif()
-mark_as_advanced(CMAKE_LLAsm_ARCHIVE)
-
-set(CMAKE_LLAsm_PREPROCESSOR_ENV_VAR "LL_PREPROCESSOR")
-set(CMAKE_LLAsm_COMPILER_ENV_VAR "LL_ASSEMBLER")
-set(CMAKE_LLAsm_ARCHIVE_ENV_VAR "LL_LINKER")
-find_file(ll_comp_in CMakeLLAsmCompiler.cmake.in PATHS ${CMAKE_ROOT}/Modules ${CMAKE_MODULE_PATH})
-# configure all variables set in this file
-configure_file(${ll_comp_in} ${CMAKE_PLATFORM_INFO_DIR}/CMakeLLAsmCompiler.cmake @ONLY)
-mark_as_advanced(ll_comp_in)
diff --git a/libclc/cmake/CMakeLLAsmCompiler.cmake.in b/libclc/cmake/CMakeLLAsmCompiler.cmake.in
deleted file mode 100644
index 2b00f69234dd22..00000000000000
--- a/libclc/cmake/CMakeLLAsmCompiler.cmake.in
+++ /dev/null
@@ -1,10 +0,0 @@
-set(CMAKE_LLAsm_PREPROCESSOR "@CMAKE_LLAsm_PREPROCESSOR@")
-set(CMAKE_LLAsm_COMPILER "@CMAKE_LLAsm_COMPILER@")
-set(CMAKE_LLAsm_ARCHIVE "@CMAKE_LLAsm_ARCHIVE@")
-set(CMAKE_LLAsm_COMPILER_LOADED 1)
-
-set(CMAKE_LLAsm_SOURCE_FILE_EXTENSIONS ll)
-set(CMAKE_LLAsm_OUTPUT_EXTENSION .bc)
-set(CMAKE_LLAsm_OUTPUT_EXTENSION_REPLACE 1)
-set(CMAKE_STATIC_LIBRARY_PREFIX_LLAsm "")
-set(CMAKE_STATIC_LIBRARY_SUFFIX_LLAsm ".bc")
diff --git a/libclc/cmake/CMakeLLAsmInformation.cmake b/libclc/cmake/CMakeLLAsmInformation.cmake
deleted file mode 100644
index 35ec3081da0f7d..00000000000000
--- a/libclc/cmake/CMakeLLAsmInformation.cmake
+++ /dev/null
@@ -1,12 +0,0 @@
-if(NOT CMAKE_LLAsm_COMPILE_OBJECT)
- set(CMAKE_LLAsm_COMPILE_OBJECT
- "${CMAKE_LLAsm_PREPROCESSOR} -E -P <DEFINES> <INCLUDES> <FLAGS> -x cl <SOURCE> -o <OBJECT>.temp"
- "<CMAKE_LLAsm_COMPILER> -o <OBJECT> <OBJECT>.temp")
-endif()
-
-if(NOT CMAKE_LLAsm_CREATE_STATIC_LIBRARY)
- set(CMAKE_LLAsm_CREATE_STATIC_LIBRARY
- "<CMAKE_LLAsm_ARCHIVE> -o <TARGET> <OBJECTS>")
-endif()
-
-set(CMAKE_INCLUDE_FLAG_LLAsm "-I")
diff --git a/libclc/cmake/CMakeTestCLCCompiler.cmake b/libclc/cmake/CMakeTestCLCCompiler.cmake
deleted file mode 100644
index 869fcc3d01ab03..00000000000000
--- a/libclc/cmake/CMakeTestCLCCompiler.cmake
+++ /dev/null
@@ -1,56 +0,0 @@
-if(CMAKE_CLC_COMPILER_FORCED)
- # The compiler configuration was forced by the user.
- # Assume the user has configured all compiler information.
- set(CMAKE_CLC_COMPILER_WORKS TRUE)
- return()
-endif()
-
-include(CMakeTestCompilerCommon)
-
-# Remove any cached result from an older CMake version.
-# We now store this in CMakeCCompiler.cmake.
-unset(CMAKE_CLC_COMPILER_WORKS CACHE)
-
-# This file is used by EnableLanguage in cmGlobalGenerator to
-# determine that that selected CLC compiler can actually compile
-# and link the most basic of programs. If not, a fatal error
-# is set and cmake stops processing commands and will not generate
-# any makefiles or projects.
-if(NOT CMAKE_CLC_COMPILER_WORKS)
- PrintTestCompilerStatus("CLC" "")
- file(WRITE ${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/CMakeTmp/testCLCCompiler.cl
- "__kernel void test_k(global int * a)\n"
- "{ *a = 1; }\n")
- try_compile(CMAKE_CLC_COMPILER_WORKS ${CMAKE_BINARY_DIR}
- ${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/CMakeTmp/testCLCCompiler.cl
- # We never generate executable so bypass the link step
- CMAKE_FLAGS -DCMAKE_CLC_LINK_EXECUTABLE='echo'
- OUTPUT_VARIABLE __CMAKE_CLC_COMPILER_OUTPUT)
- # Move result from cache to normal variable.
- set(CMAKE_CLC_COMPILER_WORKS ${CMAKE_CLC_COMPILER_WORKS})
- unset(CMAKE_CLC_COMPILER_WORKS CACHE)
- set(CLC_TEST_WAS_RUN 1)
-endif()
-
-if(NOT CMAKE_CLC_COMPILER_WORKS)
- PrintTestCompilerStatus("CLC" " -- broken")
- file(APPEND ${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/CMakeError.log
- "Determining if the CLC compiler works failed with "
- "the following output:\n${__CMAKE_CLC_COMPILER_OUTPUT}\n\n")
- message(FATAL_ERROR "The CLC compiler \"${CMAKE_CLC_COMPILER}\" "
- "is not able to compile a simple test program.\nIt fails "
- "with the following output:\n ${__CMAKE_CLC_COMPILER_OUTPUT}\n\n"
- "CMake will not be able to correctly generate this project.")
-else()
- if(CLC_TEST_WAS_RUN)
- PrintTestCompilerStatus("CLC" " -- works")
- file(APPEND ${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/CMakeOutput.log
- "Determining if the CLC compiler works passed with "
- "the following output:\n${__CMAKE_CLC_COMPILER_OUTPUT}\n\n")
- endif()
-
- include(${CMAKE_PLATFORM_INFO_DIR}/CMakeCLCCompiler.cmake)
-
-endif()
-
-unset(__CMAKE_CLC_COMPILER_OUTPUT)
diff --git a/libclc/cmake/CMakeTestLLAsmCompiler.cmake b/libclc/cmake/CMakeTestLLAsmCompiler.cmake
deleted file mode 100644
index 35948ee07a9406..00000000000000
--- a/libclc/cmake/CMakeTestLLAsmCompiler.cmake
+++ /dev/null
@@ -1,56 +0,0 @@
-if(CMAKE_LLAsm_COMPILER_FORCED)
- # The compiler configuration was forced by the user.
- # Assume the user has configured all compiler information.
- set(CMAKE_LLAsm_COMPILER_WORKS TRUE)
- return()
-endif()
-
-include(CMakeTestCompilerCommon)
-
-# Remove any cached result from an older CMake version.
-# We now store this in CMakeCCompiler.cmake.
-unset(CMAKE_LLAsm_COMPILER_WORKS CACHE)
-
-# This file is used by EnableLanguage in cmGlobalGenerator to
-# determine that that selected llvm assembler can actually compile
-# and link the most basic of programs. If not, a fatal error
-# is set and cmake stops processing commands and will not generate
-# any makefiles or projects.
-if(NOT CMAKE_LLAsm_COMPILER_WORKS)
- PrintTestCompilerStatus("LLAsm" "")
- file(WRITE ${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/CMakeTmp/testLLAsmCompiler.ll
- "define i32 @test() {\n"
- "ret i32 0 }\n" )
- try_compile(CMAKE_LLAsm_COMPILER_WORKS ${CMAKE_BINARY_DIR}
- ${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/CMakeTmp/testLLAsmCompiler.ll
- # We never generate executable so bypass the link step
- CMAKE_FLAGS -DCMAKE_LLAsm_LINK_EXECUTABLE='echo'
- OUTPUT_VARIABLE __CMAKE_LLAsm_COMPILER_OUTPUT)
- # Move result from cache to normal variable.
- set(CMAKE_LLAsm_COMPILER_WORKS ${CMAKE_LLAsm_COMPILER_WORKS})
- unset(CMAKE_LLAsm_COMPILER_WORKS CACHE)
- set(LLAsm_TEST_WAS_RUN 1)
-endif()
-
-if(NOT CMAKE_LLAsm_COMPILER_WORKS)
- PrintTestCompilerStatus("LLAsm" " -- broken")
- file(APPEND ${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/CMakeError.log
- "Determining if the LLAsm compiler works failed with "
- "the following output:\n${__CMAKE_LLAsm_COMPILER_OUTPUT}\n\n")
- message(FATAL_ERROR "The LLAsm compiler \"${CMAKE_LLAsm_COMPILER}\" "
- "is not able to compile a simple test program.\nIt fails "
- "with the following output:\n ${__CMAKE_LLAsm_COMPILER_OUTPUT}\n\n"
- "CMake will not be able to correctly generate this project.")
-else()
- if(LLAsm_TEST_WAS_RUN)
- PrintTestCompilerStatus("LLAsm" " -- works")
- file(APPEND ${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/CMakeOutput.log
- "Determining if the LLAsm compiler works passed with "
- "the following output:\n${__CMAKE_LLAsm_COMPILER_OUTPUT}\n\n")
- endif()
-
- include(${CMAKE_PLATFORM_INFO_DIR}/CMakeLLAsmCompiler.cmake)
-
-endif()
-
-unset(__CMAKE_LLAsm_COMPILER_OUTPUT)
diff --git a/libclc/cmake/modules/AddLibclc.cmake b/libclc/cmake/modules/AddLibclc.cmake
new file mode 100644
index 00000000000000..b323a4b38bf7c9
--- /dev/null
+++ b/libclc/cmake/modules/AddLibclc.cmake
@@ -0,0 +1,153 @@
+# Compiles an OpenCL C - or assembles an LL file - to bytecode
+#
+# Arguments:
+# * TRIPLE <string>
+# Target triple for which to compile the bytecode file.
+# * INPUT <string>
+# File to compile/assemble to bytecode
+# * OUTPUT <string>
+# Bytecode file to generate
+# * EXTRA_OPTS <string> ...
+# List of compiler options to use. Note that some are added by default.
+# * DEPENDENCIES <string> ...
+# List of extra dependencies to inject
+#
+# Depends on the libclc::clang and libclc::llvm-as targets for compiling and
+# assembling, respectively.
+function(compile_to_bc)
+ cmake_parse_arguments(ARG
+ ""
+ "TRIPLE;INPUT;OUTPUT"
+ "EXTRA_OPTS;DEPENDENCIES"
+ ${ARGN}
+ )
+
+ # If this is an LLVM IR file (identified soley by its file suffix),
+ # pre-process it with clang to a temp file, then assemble that to bytecode.
+ set( TMP_SUFFIX )
+ get_filename_component( FILE_EXT ${ARG_INPUT} EXT )
+ if( NOT ${FILE_EXT} STREQUAL ".ll" )
+ # Pass '-c' when not running the preprocessor
+ set( PP_OPTS -c )
+ else()
+ set( PP_OPTS -E;-P )
+ set( TMP_SUFFIX .tmp )
+ endif()
+
+ set( TARGET_ARG )
+ if( ARG_TRIPLE )
+ set( TARGET_ARG "-target" ${ARG_TRIPLE} )
+ endif()
+
+ add_custom_command(
+ OUTPUT ${ARG_OUTPUT}${TMP_SUFFIX}
+ COMMAND libclc::clang
+ ${TARGET_ARG}
+ ${PP_OPTS}
+ ${ARG_EXTRA_OPTS}
+ -MD -MF ${ARG_OUTPUT}.d -MT ${ARG_OUTPUT}${TMP_SUFFIX}
+ # LLVM 13 enables standard includes by default - we don't want
+ # those when pre-processing IR. We disable it unconditionally.
+ $<$<VERSION_GREATER_EQUAL:${LLVM_PACKAGE_VERSION},13.0.0>:-cl-no-stdinc>
+ -emit-llvm
+ -o ${ARG_OUTPUT}${TMP_SUFFIX}
+ -x cl
+ ${ARG_INPUT}
+ DEPENDS
+ libclc::clang
+ ${ARG_INPUT}
+ ${ARG_DEPENDENCIES}
+ DEPFILE ${ARG_OUTPUT}.d
+ )
+
+ if( ${FILE_EXT} STREQUAL ".ll" )
+ add_custom_command(
+ OUTPUT ${ARG_OUTPUT}
+ COMMAND libclc::llvm-as -o ${ARG_OUTPUT} ${ARG_OUTPUT}${TMP_SUFFIX}
+ DEPENDS libclc::llvm-as ${ARG_OUTPUT}${TMP_SUFFIX}
+ )
+ endif()
+endfunction()
+
+# Links together one or more bytecode files
+#
+# Arguments:
+# * OUTPUT <string>
+# Bytecode file to generate
+# * INPUT <string> ...
+# List of bytecode files to link together
+function(link_bc)
+ cmake_parse_arguments(ARG
+ ""
+ "OUTPUT"
+ "INPUTS"
+ ${ARGN}
+ )
+
+ add_custom_command(
+ OUTPUT ${ARG_OUTPUT}
+ COMMAND libclc::llvm-link -o ${ARG_OUTPUT} ${ARG_INPUTS}
+ DEPENDS libclc::llvm-link ${ARG_INPUTS}
+ )
+endfunction()
+
+# Decomposes and returns variables based on a libclc triple and architecture
+# combination. Returns data via one or more optional output variables.
+#
+# Arguments:
+# * TRIPLE <string>
+# libclc target triple to query
+# * DEVICE <string>
+# libclc device to query
+#
+# Optional Arguments:
+# * CPU <var>
+# Variable name to be set to the target CPU
+# * ARCH_SUFFIX <var>
+# Variable name to be set to the triple/architecture suffix
+# * CLANG_TRIPLE <var>
+# Variable name to be set to the normalized clang triple
+function(get_libclc_device_info)
+ cmake_parse_arguments(ARG
+ ""
+ "TRIPLE;DEVICE;CPU;ARCH_SUFFIX;CLANG_TRIPLE"
+ ""
+ ${ARGN}
+ )
+
+ if( NOT ARG_TRIPLE OR NOT ARG_DEVICE )
+ message( FATAL_ERROR "Must provide both TRIPLE and DEVICE" )
+ endif()
+
+ string( REPLACE "-" ";" TRIPLE ${ARG_TRIPLE} )
+ list( GET TRIPLE 0 ARCH )
+
+ # Some targets don't have a specific device architecture to target
+ if( ARG_DEVICE STREQUAL none OR ARCH STREQUAL spirv OR ARCH STREQUAL spirv64 )
+ set( cpu )
+ set( arch_suffix "${ARG_TRIPLE}" )
+ else()
+ set( cpu "${ARG_DEVICE}" )
+ set( arch_suffix "${ARG_DEVICE}-${ARG_TRIPLE}" )
+ endif()
+
+ if( ARG_CPU )
+ set( ${ARG_CPU} ${cpu} PARENT_SCOPE )
+ endif()
+
+ if( ARG_ARCH_SUFFIX )
+ set( ${ARG_ARCH_SUFFIX} ${arch_suffix} PARENT_SCOPE )
+ endif()
+
+ # Some libclc targets are not real clang triples: return their canonical
+ # triples.
+ if( ARCH STREQUAL spirv OR ARCH STREQUAL clspv )
+ set( t "spir--" )
+ elseif( ARCH STREQUAL spirv64 OR ARCH STREQUAL clspv64 )
+ set( t "spir64--" )
+ endif()
+
+ if( ARG_CLANG_TRIPLE )
+ set( ${ARG_CLANG_TRIPLE} ${t} PARENT_SCOPE )
+ endif()
+endfunction()
diff --git a/libclc/generic/lib/SOURCES b/libclc/generic/lib/SOURCES
index ee2736b5fbc57f..579e909e53d462 100644
--- a/libclc/generic/lib/SOURCES
+++ b/libclc/generic/lib/SOURCES
@@ -48,7 +48,6 @@ cl_khr_int64_extended_atomics/atom_max.cl
cl_khr_int64_extended_atomics/atom_min.cl
cl_khr_int64_extended_atomics/atom_or.cl
cl_khr_int64_extended_atomics/atom_xor.cl
-convert.cl
common/degrees.cl
common/mix.cl
common/radians.cl
diff --git a/llvm/tools/CMakeLists.txt b/llvm/tools/CMakeLists.txt
index cde57367934e4b..db66dad5dc0dbd 100644
--- a/llvm/tools/CMakeLists.txt
+++ b/llvm/tools/CMakeLists.txt
@@ -52,6 +52,9 @@ add_llvm_implicit_projects()
add_llvm_external_project(polly)
+# libclc depends on clang
+add_llvm_external_project(libclc)
+
# Add subprojects specified using LLVM_EXTERNAL_PROJECTS
foreach(p ${LLVM_EXTERNAL_PROJECTS})
add_llvm_external_project(${p})
>From 29c9a184c11ed0ba46a8c91b6fbeea828c9fce0d Mon Sep 17 00:00:00 2001
From: Fraser Cormack <fraser at codeplay.com>
Date: Thu, 4 Apr 2024 17:49:13 +0100
Subject: [PATCH 2/2] [libclc] Add initial LIT tests
---
libclc/CMakeLists.txt | 4 +++
libclc/test/CMakeLists.txt | 38 ++++++++++++++++++++++++++
libclc/test/add_sat.cl | 3 +++
libclc/test/as_type.cl | 5 +++-
libclc/test/convert.cl | 5 +++-
libclc/test/cos.cl | 5 +++-
libclc/test/cross.cl | 5 +++-
libclc/test/fabs.cl | 5 +++-
libclc/test/get_group_id.cl | 5 +++-
libclc/test/lit.cfg.py | 49 ++++++++++++++++++++++++++++++++++
libclc/test/lit.site.cfg.py.in | 24 +++++++++++++++++
libclc/test/rsqrt.cl | 8 ++++--
libclc/test/subsat.cl | 11 +++++---
13 files changed, 155 insertions(+), 12 deletions(-)
create mode 100644 libclc/test/CMakeLists.txt
create mode 100644 libclc/test/lit.cfg.py
create mode 100644 libclc/test/lit.site.cfg.py.in
diff --git a/libclc/CMakeLists.txt b/libclc/CMakeLists.txt
index e509a61f51b927..c3b90e2ffd66b3 100644
--- a/libclc/CMakeLists.txt
+++ b/libclc/CMakeLists.txt
@@ -403,3 +403,7 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} )
endif()
endforeach( d )
endforeach( t )
+
+if( NOT LIBCLC_STANDALONE_BUILD )
+ add_subdirectory( test )
+endif()
diff --git a/libclc/test/CMakeLists.txt b/libclc/test/CMakeLists.txt
new file mode 100644
index 00000000000000..80a981be99a02e
--- /dev/null
+++ b/libclc/test/CMakeLists.txt
@@ -0,0 +1,38 @@
+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} )
+
+ 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=${arch_suffix}.bc"
+ DEPENDS ${LIBCLC_TEST_DEPS}
+ )
+
+ add_dependencies( check-libclc check-libclc-${arch_suffix} )
+
+ endforeach()
+endforeach()
diff --git a/libclc/test/add_sat.cl b/libclc/test/add_sat.cl
index 45c8567b440397..f23a105ac83386 100644
--- a/libclc/test/add_sat.cl
+++ b/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);
}
diff --git a/libclc/test/as_type.cl b/libclc/test/as_type.cl
index e8fb1228d28d11..729303a53da051 100644
--- a/libclc/test/as_type.cl
+++ b/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);
}
diff --git a/libclc/test/convert.cl b/libclc/test/convert.cl
index 928fc326b6a18f..9ea5b192769a26 100644
--- a/libclc/test/convert.cl
+++ b/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);
}
diff --git a/libclc/test/cos.cl b/libclc/test/cos.cl
index 4230eb2a0e93c6..21384ba1d73dfd 100644
--- a/libclc/test/cos.cl
+++ b/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);
}
diff --git a/libclc/test/cross.cl b/libclc/test/cross.cl
index 08955cbd9af568..c628a6656c5739 100644
--- a/libclc/test/cross.cl
+++ b/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]);
}
diff --git a/libclc/test/fabs.cl b/libclc/test/fabs.cl
index 91d42c466676fd..27881304224567 100644
--- a/libclc/test/fabs.cl
+++ b/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);
}
diff --git a/libclc/test/get_group_id.cl b/libclc/test/get_group_id.cl
index 43725cda802710..cd4d114b59ed16 100644
--- a/libclc/test/get_group_id.cl
+++ b/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;
}
diff --git a/libclc/test/lit.cfg.py b/libclc/test/lit.cfg.py
new file mode 100644
index 00000000000000..86e69ac0ddd8e5
--- /dev/null
+++ b/libclc/test/lit.cfg.py
@@ -0,0 +1,49 @@
+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)
diff --git a/libclc/test/lit.site.cfg.py.in b/libclc/test/lit.site.cfg.py.in
new file mode 100644
index 00000000000000..b507ef5aaa240d
--- /dev/null
+++ b/libclc/test/lit.site.cfg.py.in
@@ -0,0 +1,24 @@
+ at 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"))
diff --git a/libclc/test/rsqrt.cl b/libclc/test/rsqrt.cl
index 13ad216b79f4bf..89bd903feecc91 100644
--- a/libclc/test/rsqrt.cl
+++ b/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
diff --git a/libclc/test/subsat.cl b/libclc/test/subsat.cl
index a83414b4dc850c..ec808e7556d26e 100644
--- a/libclc/test/subsat.cl
+++ b/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;
}
\ No newline at end of file
More information about the cfe-commits
mailing list