diff --git a/libclc/CMakeLists.txt b/libclc/CMakeLists.txt
index c6e3cdf23fe0c..7528228b3b7f9 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,14 @@ 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(
+       TARGET ${file}
+       INPUTS ${PROJECT_SOURCE_DIR}/generic/lib/${file}.ll
+    )
+    install( FILES $<TARGET_PROPERTY:${file},TARGET_FILE> ARCHIVE
+      DESTINATION "${CMAKE_INSTALL_DATADIR}/clc" )
+  endforeach()
 endif()
 
 find_package( Python3 REQUIRED COMPONENTS Interpreter )
@@ -232,19 +248,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 +268,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,53 +304,97 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} )
       set( opt_flags -O3 )
     endif()
 
+    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 warning
+      -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} )
-    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( 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}>
-      DEPENDS ${builtins_link_lib_tgt} )
-    add_custom_target( "opt.${obj_suffix}" ALL
-      DEPENDS ${builtins_opt_lib_tgt} )
-
-    if( ${ARCH} STREQUAL "spirv" OR ${ARCH} STREQUAL "spirv64" )
+
+    link_bc(
+      TARGET ${builtins_link_lib_tgt}
+      INPUTS ${bytecode_files}
+    )
+
+    set( builtins_link_lib $<TARGET_PROPERTY:${builtins_link_lib_tgt},TARGET_FILE> )
+
+    if( ARCH STREQUAL spirv OR ARCH STREQUAL spirv64 )
       set( spv_suffix ${arch_suffix}.spv )
-      add_custom_command( OUTPUT "${spv_suffix}"
-        COMMAND ${LLVM_SPIRV} ${spvflags} -o "${spv_suffix}" $<TARGET_FILE:${builtins_link_lib_tgt}>
-        DEPENDS ${builtins_link_lib_tgt} )
+      add_custom_command( OUTPUT ${spv_suffix}
+        COMMAND ${LLVM_SPIRV} ${spvflags} -o ${spv_suffix} ${builtins_link_lib}
+        DEPENDS ${builtins_link_lib_tgt}
+      )
       add_custom_target( "prepare-${spv_suffix}" ALL DEPENDS "${spv_suffix}" )
       install( FILES ${CMAKE_CURRENT_BINARY_DIR}/${spv_suffix}
          DESTINATION "${CMAKE_INSTALL_DATADIR}/clc" )
     else()
+      set( builtins_opt_lib_tgt builtins.opt.${arch_suffix} )
+
+      # Add opt target
+      add_custom_command( OUTPUT ${builtins_opt_lib_tgt}.bc
+        COMMAND libclc::opt ${opt_flags} -o ${builtins_opt_lib_tgt}.bc
+          ${builtins_link_lib}
+        DEPENDS libclc::opt ${builtins_link_lib_tgt}
+      )
+      add_custom_target( ${builtins_opt_lib_tgt}
+        ALL DEPENDS ${builtins_opt_lib_tgt}.bc
+      )
+      set_target_properties( ${builtins_opt_lib_tgt}
+        PROPERTIES TARGET_FILE ${builtins_opt_lib_tgt}.bc
+      )
+
       # Add prepare target
-      add_custom_command( OUTPUT "${obj_suffix}"
-        COMMAND prepare_builtins -o "${obj_suffix}" ${builtins_opt_lib_tgt}
-        DEPENDS "opt.${obj_suffix}" ${builtins_opt_lib_tgt} prepare_builtins )
-      add_custom_target( "prepare-${obj_suffix}" ALL DEPENDS "${obj_suffix}" )
+      set( obj_suffix ${arch_suffix}.bc )
+      add_custom_command( OUTPUT ${obj_suffix}
+        COMMAND prepare_builtins -o ${obj_suffix}
+          $<TARGET_PROPERTY:${builtins_opt_lib_tgt},TARGET_FILE>
+        DEPENDS ${builtins_opt_lib_tgt} prepare_builtins )
+      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,10 +402,10 @@ 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}" )
+          DEPENDS prepare-${obj_suffix} )
         install( FILES ${CMAKE_CURRENT_BINARY_DIR}/${alias_suffix} DESTINATION "${CMAKE_INSTALL_DATADIR}/clc" )
       endforeach( a )
     endif()
diff --git a/libclc/cmake/CMakeCLCCompiler.cmake.in b/libclc/cmake/CMakeCLCCompiler.cmake.in
deleted file mode 100644
index 2730b83d9e7d0..0000000000000
--- 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 95327e4439722..0000000000000
--- 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 94d85d9e666ad..0000000000000
--- 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 1c424c79cb453..0000000000000
--- 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 2b00f69234dd2..0000000000000
--- 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 35ec3081da0f7..0000000000000
--- 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 869fcc3d01ab0..0000000000000
--- 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 35948ee07a940..0000000000000
--- 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 0000000000000..5e09cde8035c2
--- /dev/null
+++ b/libclc/cmake/modules/AddLibclc.cmake
@@ -0,0 +1,156 @@
+# 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:
+# * TARGET <string>
+#     Custom target to create
+# * INPUT <string> ...
+#     List of bytecode files to link together
+function(link_bc)
+  cmake_parse_arguments(ARG
+    ""
+    "TARGET"
+    "INPUTS"
+    ${ARGN}
+  )
+
+  add_custom_command(
+    OUTPUT ${ARG_TARGET}.bc
+    COMMAND libclc::llvm-link -o ${ARG_TARGET}.bc ${ARG_INPUTS}
+    DEPENDS libclc::llvm-link ${ARG_INPUTS}
+  )
+
+  add_custom_target( ${ARG_TARGET} ALL DEPENDS ${ARG_TARGET}.bc )
+  set_target_properties( ${ARG_TARGET} PROPERTIES TARGET_FILE ${ARG_TARGET}.bc )
+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( ARG_TRIPLE "spir--" )
+  elseif( ARCH STREQUAL spirv64 OR ARCH STREQUAL clspv64 )
+    set( ARG_TRIPLE "spir64--" )
+  endif()
+
+  if( ARG_CLANG_TRIPLE )
+    set( ${ARG_CLANG_TRIPLE} ${ARG_TRIPLE} PARENT_SCOPE )
+  endif()
+endfunction()
diff --git a/libclc/generic/lib/SOURCES b/libclc/generic/lib/SOURCES
index ee2736b5fbc57..579e909e53d46 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 cde57367934e4..db66dad5dc0db 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})