[libc] Add support for creating wrapper headers for offloading in clang

[jhuber6]

Summary

This is an alternate approach to the patches proposed in D153897 and
D153794. Rather than exporting a single header that can be included on
the GPU in all circumstances, this patch chooses to instead generate a
separate set of headers that only provides the declarations. This can
then be used by external tooling to set up what's on the GPU. This
leaves room for header hacks for offloading languages without needing to
worry about the libc implementation.

Currently this generates a set of headers that only contain the
declarations. These will then be installed to a new clang resource
directory called llvm_libc_wrappers/ which will house the shim code.
We can then automaticlaly include this from clang when offloading to
wrap around the headers while specifying what's on the GPU.

[jhuber6]

For reference, here is what one of the newly generated headers looks like that is used.

#ifndef __LLVM_LIBC_DECLARATIONS_STDIO_H
#define __LLVM_LIBC_DECLARATIONS_STDIO_H

#ifndef __LIBC_ATTRS
#define __LIBC_ATTRS
#endif

#ifdef __cplusplus
extern "C" {
#endif

int puts(const char *__restrict) __LIBC_ATTRS;

int fputs(const char *__restrict, FILE *__restrict) __LIBC_ATTRS;

extern FILE * stdin __LIBC_ATTRS;
extern FILE * stdout __LIBC_ATTRS;
extern FILE * stderr __LIBC_ATTRS;

#ifdef __cplusplus
}
#endif

Unfortunately I have already run into a few problems with the re-declarations given the GNU libc headers. Here is the error message when including string.h now,

/home/jhuber/Documents/llvm/clang/lib/clang/17/include/llvm_libc_wrappers/llvm-libc-decls/string.h:54:8: error: 'strstr' is missing exception specification 'noexcept(true)'
   54 | char * strstr(const char *, const char *) __LIBC_ATTRS;
      |        ^
/usr/include/string.h:343:1: note: previous declaration is here
  343 | strstr (const char *__haystack, const char *__needle) __THROW
      | ^
5 errors generated.

This occurs for memchr, strchr, strpbrk, strchr, and strstr. If you define __LIBC_ATTRS to the noexcept(true) you get a different error,

/home/jhuber/Documents/llvm/clang/lib/clang/17/include/llvm_libc_wrappers/llvm-libc-decls/string.h:54:8: error: functions that differ only in their return type cannot be overloaded
   54 | char * strstr(const char *, const char *) __LIBC_ATTRS;
      | ~~~~~~ ^
/usr/include/string.h:343:1: note: previous definition is here
  343 | __extern_always_inline const char *
      |                              ~~~~~~
  344 | strstr (const char *__haystack, const char *__needle) __THROW
      | ^

Looking at the definitions, they look like this in the GNU headers,

extern char *strstr (char *__haystack, const char *__needle)    
     __THROW __asm ("strstr") __attribute_pure__ __nonnull ((1, 2));    
extern const char *strstr (const char *__haystack, const char *__needle)    
     __THROW __asm ("strstr") __attribute_pure__ __nonnull ((1, 2));

Does anyone have any suggestions on working around this? The other supported headers work as far as I can tell.

[jhuber6]

Hack around the string problem. GNU likes to provide different prototypes for C++. Manually disable this for now. Unsure if this will have reasonable fallout, but it seems bizarre that string.h would define C++ constructs?

[jhuber6]

Semi-fix hack for string.h and fix ctype.h. string.h required undefining
C++ mode so we didn't use weird GNU C++ handling, which we then still need the
extern "C" for. The cytpe problems come from GNU defining everything as a
macro so it fails to redeclare.

The amount of hacks that just this has required so far is fairly convincing to
me that this is the more correct solution and should be separate from libc.

[jhuber6]

Add some checks to stdlib.h to ensure ABI compatibility for div functions.

Because this patch makes us always include the LLVM libc repo when offloading,
we mask off the wrappers if they were not installed by libc, so this is simply
a passthrough include header if the user didn't build with LLVM C library for
the GPU.

[jhuber6]

Fix guard on the headers for offloading languages

[sivachandra]

The libc specific changes are really minimal straightforward. The GPU side, the clang driver changes etc. need a review by a GPU expert.

libc/include/CMakeLists.txt

8

Where does this come from?

[jhuber6]

libc/include/CMakeLists.txt

8

It's a global CMake module that LLVM provides in cmake/Modules/GetClangResourceDir.cmake. I only expect the GPU build to be done in-tree.

[JonChesterfield]

OK, let's go with this. It's a fairly alarming mess localised quite closely to the language that requires the complexity, minimal damage to libc itself.

[jdoerfert]

If it compiles, ship it.

[jhuber6]

Changing this to only apply to OpenMP for now. It breaks CUDA / HIP builds
because they already have forward declarations of things like malloc or
memcpy on the GPU that conflict. We'll need to clean those up later.

clang/lib/Driver/ToolChains/Clang.cpp

void Clang::AddPreprocessingOptions(Compilation &C, const JobAction &JA,
   // from e.g. /usr/local/include.
   if (JA.isOffloading(Action::OFK_Cuda))
     getToolChain().AddCudaIncludeArgs(Args, CmdArgs);
   if (JA.isOffloading(Action::OFK_HIP))
     getToolChain().AddHIPIncludeArgs(Args, CmdArgs);
 
   // If we are compiling for a GPU target we want to override the system headers
   // with ones created by the 'libc' project if present.
-  // FIXME: We need to find a way to make these headers compatible with the
-  // host environment so they can be included from offloading languages. For now
-  // these are only active when targeting the GPU with cross-compilation.
   if (!Args.hasArg(options::OPT_nostdinc) &&
       !Args.hasArg(options::OPT_nogpuinc) &&
       !Args.hasArg(options::OPT_nobuiltininc) &&
-      C.getActiveOffloadKinds() == Action::OFK_None &&
       (getToolChain().getTriple().isNVPTX() ||
        getToolChain().getTriple().isAMDGCN())) {
 
-      // Add include/gpu-none-libc/* to our system include path. This lets us use
-      // GPU-specific system headers first.
+    // Without an offloading language we will include these headers directly.
+    // Offloading languages will instead only use the declarations stored in
+    // the resource directory at clang/lib/Headers/llvm_libc_wrappers.
+    if (C.getActiveOffloadKinds() == Action::OFK_None) {
       SmallString<128> P(llvm::sys::path::parent_path(D.InstalledDir));
       llvm::sys::path::append(P, "include");
       llvm::sys::path::append(P, "gpu-none-llvm");
       CmdArgs.push_back("-c-isystem");
       CmdArgs.push_back(Args.MakeArgString(P));
+    } else if (C.getActiveOffloadKinds() == Action::OFK_OpenMP) {
+      // TODO: CUDA / HIP include their own headers for some common functions
+      // implemented here. We'll need to clean those up so they do not conflict.
+      SmallString<128> P(D.ResourceDir);
+      llvm::sys::path::append(P, "include");
+      llvm::sys::path::append(P, "llvm_libc_wrappers");
+      CmdArgs.push_back("-internal-isystem");
+      CmdArgs.push_back(Args.MakeArgString(P));
+    }
   }
 
   // If we are offloading to a target via OpenMP we need to include the
   // openmp_wrappers folder which contains alternative system headers.
   if (JA.isDeviceOffloading(Action::OFK_OpenMP) &&
       !Args.hasArg(options::OPT_nostdinc) &&
       !Args.hasArg(options::OPT_nogpuinc) &&
       (getToolChain().getTriple().isNVPTX() ||

clang/lib/Headers/CMakeLists.txt

set(openmp_wrapper_files
   openmp_wrappers/cmath
   openmp_wrappers/complex.h
   openmp_wrappers/complex
   openmp_wrappers/__clang_openmp_device_functions.h
   openmp_wrappers/complex_cmath.h
   openmp_wrappers/new
 )
 
+set(llvm_libc_wrapper_files
+  llvm_libc_wrappers/stdio.h
+  llvm_libc_wrappers/stdlib.h
+  llvm_libc_wrappers/string.h
+  llvm_libc_wrappers/ctype.h
+  llvm_libc_wrappers/inttypes.h
+)
+
 include(GetClangResourceDir)
 get_clang_resource_dir(output_dir PREFIX ${LLVM_LIBRARY_OUTPUT_INTDIR}/.. SUBDIR include)
 set(out_files)
 set(generated_files)
 
 set(arm_common_generated_files)
 set(arm_only_generated_files)
 set(aarch64_only_generated_files)
function(clang_generate_header td_option td_file out_file)
   set(out_files ${out_files} PARENT_SCOPE)
   list(APPEND generated_files "${CMAKE_CURRENT_BINARY_DIR}/${out_file}")
   set(generated_files ${generated_files} PARENT_SCOPE)
 endfunction(clang_generate_header)
 
 
 # Copy header files from the source directory to the build directory
 foreach( f ${files} ${cuda_wrapper_files} ${cuda_wrapper_bits_files}
-           ${ppc_wrapper_files} ${openmp_wrapper_files} ${hlsl_files})
+           ${ppc_wrapper_files} ${openmp_wrapper_files} ${hlsl_files}
+           ${llvm_libc_wrapper_files})
   copy_header_to_output_dir(${CMAKE_CURRENT_SOURCE_DIR} ${f})
 endforeach( f )
 
 # Generate header files and copy them to the build directory
 if(ARM IN_LIST LLVM_TARGETS_TO_BUILD OR AArch64 IN_LIST LLVM_TARGETS_TO_BUILD)
   # Generate arm_neon.h
   clang_generate_header(-gen-arm-neon arm_neon.td arm_neon.h)
   # Generate arm_fp16.h
add_dependencies("clang-resource-headers"
                  "ppc-htm-resource-headers"
                  "riscv-resource-headers"
                  "systemz-resource-headers"
                  "ve-resource-headers"
                  "webassembly-resource-headers"
                  "x86-resource-headers"
                  "opencl-resource-headers"
                  "openmp-resource-headers"
+                 "llvm-libc-resource-headers"
                  "windows-resource-headers"
                  "utility-resource-headers")
 
 # Core/common headers
 add_header_target("core-resource-headers" ${core_files})
 add_header_target("arm-common-resource-headers" "${arm_common_files};${arm_common_generated_files}")
 
 # Architecture/platform specific targets

 add_header_target("systemz-resource-headers" "${systemz_files}")
 add_header_target("ve-resource-headers" "${ve_files}")
 add_header_target("webassembly-resource-headers" "${webassembly_files}")
 add_header_target("x86-resource-headers" "${x86_files}")
 
 # Other header groupings
 add_header_target("hlsl-resource-headers" ${hlsl_files})
 add_header_target("opencl-resource-headers" ${opencl_files})
+add_header_target("llvm-libc-resource-headers" ${llvm_libc_wrapper_files})
 add_header_target("openmp-resource-headers" ${openmp_wrapper_files})
 add_header_target("windows-resource-headers" ${windows_only_files})
 add_header_target("utility-resource-headers" ${utility_files})
 
 get_clang_resource_dir(header_install_dir SUBDIR include)
 
 #############################################################
 # Install rules for the catch-all clang-resource-headers target
install(
   COMPONENT clang-resource-headers)
 
 install(
   FILES ${ppc_wrapper_files}
   DESTINATION ${header_install_dir}/ppc_wrappers
   COMPONENT clang-resource-headers)
 
 install(
+  FILES ${llvm_libc_wrapper_files}
+  DESTINATION ${header_install_dir}/llvm_libc_wrappers
+  COMPONENT clang-resource-headers)
+
+install(
   FILES ${openmp_wrapper_files}
   DESTINATION ${header_install_dir}/openmp_wrappers
   COMPONENT clang-resource-headers)
 
 #############################################################
 # Install rules for separate header lists
 install(
   FILES ${core_files}

 
 install(
   FILES ${openmp_wrapper_files}
   DESTINATION ${header_install_dir}/openmp_wrappers
   EXCLUDE_FROM_ALL
   COMPONENT openmp-resource-headers)
 
 install(
+  FILES ${openmp_wrapper_files}
+  DESTINATION ${header_install_dir}/openmp_wrappers
+  EXCLUDE_FROM_ALL
+  COMPONENT openmp-resource-headers)
+
+install(
   FILES ${utility_files}
   DESTINATION ${header_install_dir}
   EXCLUDE_FROM_ALL
   COMPONENT utility-resource-headers)
 
 install(
   FILES ${windows_only_files}
   DESTINATION ${header_install_dir}

clang/lib/Headers/llvm_libc_wrappers/ctype.h

+//===-- Wrapper for C standard ctype.h declarations on the GPU ------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CLANG_LLVM_LIBC_WRAPPERS_CTYPE_H__
+#define __CLANG_LLVM_LIBC_WRAPPERS_CTYPE_H__
+
+#if !defined(_OPENMP) && !defined(__HIP__) && !defined(__CUDA__)
+#error "This file is for GPU offloading compilation only"
+#endif
+
+#include_next <ctype.h>
+
+#if __has_include(<llvm-libc-decls/ctype.h>)
+
+#if defined(__HIP__) || defined(__CUDA__)
+#define __LIBC_ATTRS __attribute__((device))
+#endif
+
+// The GNU headers like to provide these as macros, we need to undefine them so
+// they do not conflict with the following definitions for the GPU.
+
+#pragma push_macro("isalnum")
+#pragma push_macro("isalpha")
+#pragma push_macro("isblank")
+#pragma push_macro("iscntrl")
+#pragma push_macro("isdigit")
+#pragma push_macro("isgraph")
+#pragma push_macro("islower")
+#pragma push_macro("isprint")
+#pragma push_macro("ispunct")
+#pragma push_macro("isspace")
+#pragma push_macro("isupper")
+#pragma push_macro("isxdigit")
+#pragma push_macro("tolower")
+#pragma push_macro("toupper")
+
+#undef isalnum
+#undef isalpha
+#undef iscntrl
+#undef isdigit
+#undef islower
+#undef isgraph
+#undef isprint
+#undef ispunct
+#undef isspace
+#undef isupper
+#undef isblank
+#undef isxdigit
+#undef tolower
+#undef toupper
+
+#pragma omp begin declare target
+
+#include <llvm-libc-decls/ctype.h>
+
+#pragma omp end declare target
+
+// Restore the original macros when compiling on the host.
+#if !defined(__NVPTX__) && !defined(__AMDGPU__)
+#pragma pop_macro("isalnum")
+#pragma pop_macro("isalpha")
+#pragma pop_macro("isblank")
+#pragma pop_macro("iscntrl")
+#pragma pop_macro("isdigit")
+#pragma pop_macro("isgraph")
+#pragma pop_macro("islower")
+#pragma pop_macro("isprint")
+#pragma pop_macro("ispunct")
+#pragma pop_macro("isspace")
+#pragma pop_macro("isupper")
+#pragma pop_macro("isxdigit")
+#pragma pop_macro("tolower")
+#pragma pop_macro("toupper")
+#endif
+
+#undef __LIBC_ATTRS
+
+#endif
+
+#endif // __CLANG_LLVM_LIBC_WRAPPERS_CTYPE_H__

clang/lib/Headers/llvm_libc_wrappers/inttypes.h

+//===-- Wrapper for C standard inttypes.h declarations on the GPU ---------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CLANG_LLVM_LIBC_WRAPPERS_INTTYPES_H__
+#define __CLANG_LLVM_LIBC_WRAPPERS_INTTYPES_H__
+
+#if !defined(_OPENMP) && !defined(__HIP__) && !defined(__CUDA__)
+#error "This file is for GPU offloading compilation only"
+#endif
+
+#include_next <inttypes.h>
+
+#if __has_include(<llvm-libc-decls/inttypes.h>)
+
+#if defined(__HIP__) || defined(__CUDA__)
+#define __LIBC_ATTRS __attribute__((device))
+#endif
+
+#pragma omp begin declare target
+
+#include <llvm-libc-decls/inttypes.h>
+
+#pragma omp end declare target
+
+#undef __LIBC_ATTRS
+
+#endif
+
+#endif // __CLANG_LLVM_LIBC_WRAPPERS_INTTYPES_H__

clang/lib/Headers/llvm_libc_wrappers/llvm-libc-decls/README.txt

+LLVM libc declarations
+======================
+
+This directory will be filled by the `libc` project with declarations that are
+availible on the device. Each declaration will use the `__LIBC_ATTRS` attribute
+to control emission on the device side.

clang/lib/Headers/llvm_libc_wrappers/stdio.h

+//===-- Wrapper for C standard stdio.h declarations on the GPU ------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CLANG_LLVM_LIBC_WRAPPERS_STDIO_H__
+#define __CLANG_LLVM_LIBC_WRAPPERS_STDIO_H__
+
+#if !defined(_OPENMP) && !defined(__HIP__) && !defined(__CUDA__)
+#error "This file is for GPU offloading compilation only"
+#endif
+
+#include_next <stdio.h>
+
+#if __has_include(<llvm-libc-decls/stdio.h>)
+
+#if defined(__HIP__) || defined(__CUDA__)
+#define __LIBC_ATTRS __attribute__((device))
+#endif
+
+#pragma omp begin declare target
+
+#include <llvm-libc-decls/stdio.h>
+
+#pragma omp end declare target
+
+#undef __LIBC_ATTRS
+
+#endif
+
+#endif // __CLANG_LLVM_LIBC_WRAPPERS_STDIO_H__

clang/lib/Headers/llvm_libc_wrappers/stdlib.h

+//===-- Wrapper for C standard stdlib.h declarations on the GPU -----------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CLANG_LLVM_LIBC_WRAPPERS_STDLIB_H__
+#define __CLANG_LLVM_LIBC_WRAPPERS_STDLIB_H__
+
+#if !defined(_OPENMP) && !defined(__HIP__) && !defined(__CUDA__)
+#error "This file is for GPU offloading compilation only"
+#endif
+
+#include_next <stdlib.h>
+
+#if __has_include(<llvm-libc-decls/stdlib.h>)
+
+#if defined(__HIP__) || defined(__CUDA__)
+#define __LIBC_ATTRS __attribute__((device))
+#endif
+
+#pragma omp begin declare target
+
+// The LLVM C library uses this type so we forward declare it.
+typedef void (*__atexithandler_t)(void);
+
+// Enforce ABI compatibility with the structs used by the LLVM C library.
+_Static_assert(__builtin_offsetof(div_t, quot) == 0, "ABI mismatch!");
+_Static_assert(__builtin_offsetof(ldiv_t, quot) == 0, "ABI mismatch!");
+_Static_assert(__builtin_offsetof(lldiv_t, quot) == 0, "ABI mismatch!");
+
+#include <llvm-libc-decls/stdlib.h>
+
+#pragma omp end declare target
+
+#undef __LIBC_ATTRS
+
+#endif
+
+#endif // __CLANG_LLVM_LIBC_WRAPPERS_STDLIB_H__

clang/lib/Headers/llvm_libc_wrappers/string.h

+//===-- Wrapper for C standard string.h declarations on the GPU -----------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CLANG_LLVM_LIBC_WRAPPERS_STRING_H__
+#define __CLANG_LLVM_LIBC_WRAPPERS_STRING_H__
+
+#if !defined(_OPENMP) && !defined(__HIP__) && !defined(__CUDA__)
+#error "This file is for GPU offloading compilation only"
+#endif
+
+// The GNU headers provide non C-standard headers when in C++ mode. Manually
+// undefine it here so that the definitions agree with the C standard for our
+// purposes.
+#ifdef __cplusplus
+extern "C" {
+#pragma push_macro("__cplusplus")
+#undef __cplusplus
+#endif
+
+#include_next <string.h>
+
+#pragma pop_macro("__cplusplus")
+#ifdef __cplusplus
+}
+#endif
+
+#if __has_include(<llvm-libc-decls/string.h>)
+
+#if defined(__HIP__) || defined(__CUDA__)
+#define __LIBC_ATTRS __attribute__((device))
+#endif
+
+#pragma omp begin declare target
+
+#include <llvm-libc-decls/string.h>
+
+#pragma omp end declare target
+
+#undef __LIBC_ATTRS
+
+#endif
+
+#endif // __CLANG_LLVM_LIBC_WRAPPERS_STRING_H__

clang/test/Driver/gpu-libc-headers.c

 // REQUIRES: nvptx-registered-target
 // REQUIRES: amdgpu-registered-target
 
+// RUN:   %clang -### --target=x86_64-unknown-linux-gnu -fopenmp=libomp --sysroot=./ \
+// RUN:     -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa --offload-arch=gfx908  \
+// RUN:     -nogpulib %s 2>&1 | FileCheck %s --check-prefix=CHECK-HEADERS
+// RUN:   %clang -### --target=x86_64-unknown-linux-gnu -fopenmp=libomp --sysroot=./ \
+// RUN:     -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target=nvptx64-nvidia-cuda --offload-arch=sm_70  \
+// RUN:     -nogpulib %s 2>&1 | FileCheck %s --check-prefix=CHECK-HEADERS
+// CHECK-HEADERS: "-cc1"{{.*}}"-internal-isystem" "{{.*}}include{{.*}}llvm_libc_wrappers"{{.*}}"-isysroot" "./"
+
 // RUN:   %clang -### --target=amdgcn-amd-amdhsa -mcpu=gfx1030 -nogpulib \
 // RUN:     -nogpuinc %s 2>&1 | FileCheck %s --check-prefix=CHECK-HEADERS-DISABLED
 // RUN:   %clang -### --target=amdgcn-amd-amdhsa -mcpu=gfx1030 -nogpulib \
 // RUN:     -nostdinc %s 2>&1 | FileCheck %s --check-prefix=CHECK-HEADERS-DISABLED
 // RUN:   %clang -### --target=amdgcn-amd-amdhsa -mcpu=gfx1030 -nogpulib \
 // RUN:     -nobuiltininc %s 2>&1 | FileCheck %s --check-prefix=CHECK-HEADERS-DISABLED
 // CHECK-HEADERS-DISABLED-NOT: "-cc1"{{.*}}"-c-isystem" "{{.*}}include{{.*}}gpu-none-llvm"

libc/cmake/modules/LLVMLibCHeaderRules.cmake

COMMAND ${hdrgen_exe} -o ${out_file} --header ${ADD_GEN_HDR_GEN_HDR}
            ${LIBC_SOURCE_DIR}/config/${LIBC_TARGET_OS}/api.td
 
     WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}
     DEPENDS ${in_file} ${fq_data_files} ${td_includes}
             ${LIBC_SOURCE_DIR}/config/${LIBC_TARGET_OS}/api.td
             ${hdrgen_deps}
   )
 
+  if(LIBC_TARGET_ARCHITECTURE_IS_GPU)
+    file(MAKE_DIRECTORY ${LIBC_INCLUDE_DIR}/llvm-libc-decls)
+    set(decl_out_file ${LIBC_INCLUDE_DIR}/llvm-libc-decls/${relative_path})
+    add_custom_command(
+      OUTPUT ${decl_out_file}
+      COMMAND ${hdrgen_exe} -o ${decl_out_file}
+              --header ${ADD_GEN_HDR_GEN_HDR} --def ${in_file} --export-decls
+              ${replacement_params} -I ${LIBC_SOURCE_DIR} ${ENTRYPOINT_NAME_LIST_ARG}
+              ${LIBC_SOURCE_DIR}/config/${LIBC_TARGET_OS}/api.td
+
+      WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}
+      DEPENDS ${in_file} ${fq_data_files} ${td_includes}
+              ${LIBC_SOURCE_DIR}/config/${LIBC_TARGET_OS}/api.td
+              ${hdrgen_deps}
+    )
+  endif()
+
   if(ADD_GEN_HDR_DEPENDS)
     get_fq_deps_list(fq_deps_list ${ADD_GEN_HDR_DEPENDS})
     # Dependencies of a add_header target can only be another add_gen_header target
     # or an add_header target.
     foreach(dep IN LISTS fq_deps_list)
       get_target_property(header_file ${dep} HEADER_FILE_PATH)
       if(NOT header_file)
         message(FATAL_ERROR "Invalid dependency '${dep}' for '${fq_target_name}'.")
       endif()
     endforeach()
   endif()
   add_custom_target(
     ${fq_target_name}
-    DEPENDS ${out_file} ${fq_deps_list}
+    DEPENDS ${out_file} ${fq_deps_list} ${decl_out_file}
   )
 
   set_target_properties(
     ${fq_target_name}
     PROPERTIES
       HEADER_FILE_PATH ${out_file}
+      DECLS_FILE_PATH ${decl_out_file}
       DEPS "${fq_deps_list}"
   )
 endfunction(add_gen_header)

libc/include/CMakeLists.txt

 set(LIBC_INCLUDE_SOURCE_DIR ${CMAKE_CURRENT_SOURCE_DIR})
 set(LIBC_INCLUDE_BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR})
 
 include(LLVMLibCHeaderRules)
 
+# The GPU build wants to install files in the compiler's resource directory.
+if(LIBC_TARGET_ARCHITECTURE_IS_GPU)
+  include(GetClangResourceDir)
+endif()
+
 add_subdirectory(llvm-libc-macros)
 add_subdirectory(llvm-libc-types)
 
 add_header(
   llvm_libc_common_h
   HDR
     __llvm-libc-common.h
 )
if(NOT header_file)
     message(FATAL_ERROR "Installable header file '${target}' does not have the "
                         "HEADER_FILE_PATH property set.")
   endif()
   file(RELATIVE_PATH relative_path ${LIBC_INCLUDE_BINARY_DIR} ${header_file})
   get_filename_component(nested_dir ${relative_path} DIRECTORY)
   install(FILES ${header_file}
           DESTINATION ${LIBC_INSTALL_INCLUDE_DIR}/${nested_dir}
           COMPONENT libc-headers)
+  # The GPU optionally provides the supported declarations externally so
+  # offloading languages like CUDA and OpenMP know what is supported by libc. We
+  # install these in the compiler's resource directory at a preset location.
+  if(LIBC_TARGET_ARCHITECTURE_IS_GPU)
+    get_target_property(decls_file ${target} DECLS_FILE_PATH)
+    if(NOT decls_file)
+      continue()
+    endif()
+    get_clang_resource_dir(resource_dir SUBDIR include)
+    file(RELATIVE_PATH relative_path ${LIBC_INCLUDE_BINARY_DIR} ${decls_file})
+    get_filename_component(nested_dir ${relative_path} DIRECTORY)
+    set(install_dir
+        ${CMAKE_INSTALL_PREFIX}/${resource_dir}/llvm_libc_wrappers/${nested_dir})
+    install(FILES ${decls_file}
+            DESTINATION ${install_dir}
+            COMPONENT libc-headers)
+  endif()
 endforeach()

libc/utils/HdrGen/Generator.h

 public:
   Generator(const std::string &DefFile, const std::vector<std::string> &EN,
             const std::string &Header,
             std::unordered_map<std::string, std::string> &Map)
       : HeaderDefFile(DefFile), EntrypointNameList(EN), StdHeader(Header),
         ArgMap(Map) {}
 
   void generate(llvm::raw_ostream &OS, llvm::RecordKeeper &Records);
+  void generateDecls(llvm::raw_ostream &OS, llvm::RecordKeeper &Records);
 };
 
 } // namespace llvm_libc
 
 #endif // LLVM_LIBC_UTILS_HDRGEN_GENERATOR_H

libc/utils/HdrGen/Generator.cpp

 //===-- Implementation of the main header generation class ----------------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 
 #include "Generator.h"
 
 #include "IncludeFileCommand.h"
 #include "PublicAPICommand.h"
+#include "utils/LibcTableGenUtil/APIIndexer.h"
 
 #include "llvm/ADT/StringRef.h"
 #include "llvm/Support/MemoryBuffer.h"
 #include "llvm/Support/SourceMgr.h"
 #include "llvm/Support/raw_ostream.h"
 
 #include <cstdlib>
 #include <memory>
if (Line.startswith(CommandPrefix)) {
       OS << P.first << "\n";
     }
 
     if (P.second.empty())
       break;
   }
 }
 
+void Generator::generateDecls(llvm::raw_ostream &OS,
+                              llvm::RecordKeeper &Records) {
+
+  OS << "//===-- C standard declarations for " << StdHeader << " "
+     << std::string(80 - (42 + StdHeader.size()), '-') << "===//\n"
+     << "//\n"
+     << "// Part of the LLVM Project, under the Apache License v2.0 with LLVM "
+        "Exceptions.\n"
+     << "// See https://llvm.org/LICENSE.txt for license information.\n"
+     << "// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception\n"
+     << "//\n"
+     << "//"
+        "===-------------------------------------------------------------------"
+        "---===//\n\n";
+
+  std::string HeaderGuard(StdHeader.size(), '\0');
+  llvm::transform(StdHeader, HeaderGuard.begin(), [](const char C) -> char {
+    return !isalnum(C) ? '_' : llvm::toUpper(C);
+  });
+  OS << "#ifndef __LLVM_LIBC_DECLARATIONS_" << HeaderGuard << "\n"
+     << "#define __LLVM_LIBC_DECLARATIONS_" << HeaderGuard << "\n\n";
+
+  OS << "#ifndef __LIBC_ATTRS\n"
+     << "#define __LIBC_ATTRS\n"
+     << "#endif\n\n";
+
+  OS << "#ifdef __cplusplus\n"
+     << "extern \"C\" {\n"
+     << "#endif\n\n";
+
+  APIIndexer G(StdHeader, Records);
+  for (auto &Name : EntrypointNameList) {
+    // Filter out functions not exported by this header.
+    if (G.FunctionSpecMap.find(Name) == G.FunctionSpecMap.end())
+      continue;
+
+    llvm::Record *FunctionSpec = G.FunctionSpecMap[Name];
+    llvm::Record *RetValSpec = FunctionSpec->getValueAsDef("Return");
+    llvm::Record *ReturnType = RetValSpec->getValueAsDef("ReturnType");
+
+    OS << G.getTypeAsString(ReturnType) << " " << Name << "(";
+
+    auto ArgsList = FunctionSpec->getValueAsListOfDefs("Args");
+    for (size_t i = 0; i < ArgsList.size(); ++i) {
+      llvm::Record *ArgType = ArgsList[i]->getValueAsDef("ArgType");
+      OS << G.getTypeAsString(ArgType);
+      if (i < ArgsList.size() - 1)
+        OS << ", ";
+    }
+
+    OS << ") __LIBC_ATTRS;\n\n";
+  }
+
+  // Make another pass over entrypoints to emit object declarations.
+  for (const auto &Name : EntrypointNameList) {
+    if (G.ObjectSpecMap.find(Name) == G.ObjectSpecMap.end())
+      continue;
+    llvm::Record *ObjectSpec = G.ObjectSpecMap[Name];
+    auto Type = ObjectSpec->getValueAsString("Type");
+    OS << "extern " << Type << " " << Name << " __LIBC_ATTRS;\n";
+  }
+
+  // Emit a final newline if we emitted any object declarations.
+  if (llvm::any_of(EntrypointNameList, [&](const std::string &Name) {
+        return G.ObjectSpecMap.find(Name) != G.ObjectSpecMap.end();
+      }))
+    OS << "\n";
+
+  OS << "#ifdef __cplusplus\n"
+     << "}\n"
+     << "#endif\n\n";
+  OS << "#endif\n";
+}
+
 } // namespace llvm_libc

libc/utils/HdrGen/Main.cpp

 llvm::cl::list<std::string> EntrypointNamesOption(
     "e", llvm::cl::value_desc("<list of entrypoints>"),
     llvm::cl::desc(
         "Each --e is one entrypoint (generated from entrypoints.txt)"),
     llvm::cl::OneOrMore);
 llvm::cl::list<std::string> ReplacementValues(
     "args", llvm::cl::desc("Command separated <argument name>=<value> pairs."),
     llvm::cl::value_desc("<name=value>[,name=value]"));
+llvm::cl::opt<bool> ExportDecls(
+    "export-decls",
+    llvm::cl::desc("Output a new header containing only the entrypoints."));
 
 void ParseArgValuePairs(std::unordered_map<std::string, std::string> &Map) {
   for (std::string &R : ReplacementValues) {
     auto Pair = llvm::StringRef(R).split('=');
     Map[std::string(Pair.first)] = std::string(Pair.second);
   }
 }
 
 } // anonymous namespace
 
 namespace llvm_libc {
 
 bool HeaderGeneratorMain(llvm::raw_ostream &OS, llvm::RecordKeeper &Records) {
   std::unordered_map<std::string, std::string> ArgMap;
   ParseArgValuePairs(ArgMap);
   Generator G(HeaderDefFile, EntrypointNamesOption, StandardHeader, ArgMap);
+  if (ExportDecls)
+    G.generateDecls(OS, Records);
+  else
     G.generate(OS, Records);
 
   return false;
 }
 
 } // namespace llvm_libc
 
 int main(int argc, char *argv[]) {
   llvm::cl::ParseCommandLineOptions(argc, argv);
   return TableGenMain(argv[0], &llvm_libc::HeaderGeneratorMain);
 }