Skip to content

Conversation

@jhuber6
Copy link
Contributor

@jhuber6 jhuber6 commented Nov 17, 2025

Summary:
These were originally intended to represent the functions that are
present on the GPU as to be provided by the LLVM libc implementation.
The original plan was that LLVM libc would report which functions were
supported and then the offload interface would mark those as supported.
The problem is that these wrapper headers are very difficult to make
work given the various libc extensions everyone does so they were
extremely fragile.

OpenMP already declares all functions used inside of a target region as
implicitly host / device, while these headers weren't even used for CUDA
/ HIP yet anyway. The only things we need to define right now are the
stdio FILE types. If we want to make this work for CUDA we'd need to
define these manually, but we're a ways off and that's way easier
because they do proper overloading.

Summary:
These were originally intended to represent the functions that are
present on the GPU as to be provided by the LLVM libc implementation.
The original plan was that LLVM libc would report which functions were
supported and then the offload interface would mark those as supported.
The problem is that these wrapper headers are very difficult to make
work given the various libc extensions everyone does so they were
extremely fragile.

OpenMP already declares all functions used inside of a target region as
implicitly host / device, while these headers weren't even used for CUDA
/ HIP yet anyway. The only things we need to define right now are the
stdio FILE types. If we want to make this work for CUDA we'd need to
define these manually, but we're a ways off and that's way easier
because they do proper overloading.
@llvmbot llvmbot added clang Clang issues not falling into any other category backend:X86 clang:headers Headers provided by Clang, e.g. for intrinsics libc labels Nov 17, 2025
@llvmbot
Copy link
Member

llvmbot commented Nov 17, 2025

@llvm/pr-subscribers-backend-x86

@llvm/pr-subscribers-libc

Author: Joseph Huber (jhuber6)

Changes

Summary:
These were originally intended to represent the functions that are
present on the GPU as to be provided by the LLVM libc implementation.
The original plan was that LLVM libc would report which functions were
supported and then the offload interface would mark those as supported.
The problem is that these wrapper headers are very difficult to make
work given the various libc extensions everyone does so they were
extremely fragile.

OpenMP already declares all functions used inside of a target region as
implicitly host / device, while these headers weren't even used for CUDA
/ HIP yet anyway. The only things we need to define right now are the
stdio FILE types. If we want to make this work for CUDA we'd need to
define these manually, but we're a ways off and that's way easier
because they do proper overloading.


Full diff: https://github.com/llvm/llvm-project/pull/168438.diff

12 Files Affected:

  • (modified) clang/lib/Headers/llvm_libc_wrappers/assert.h (+3-5)
  • (modified) clang/lib/Headers/llvm_libc_wrappers/ctype.h (+3-115)
  • (modified) clang/lib/Headers/llvm_libc_wrappers/inttypes.h (+3-5)
  • (removed) clang/lib/Headers/llvm_libc_wrappers/llvm-libc-decls/README.txt (-6)
  • (modified) clang/lib/Headers/llvm_libc_wrappers/stdio.h (+10-38)
  • (modified) clang/lib/Headers/llvm_libc_wrappers/stdlib.h (+3-24)
  • (modified) clang/lib/Headers/llvm_libc_wrappers/string.h (+2-70)
  • (modified) clang/lib/Headers/llvm_libc_wrappers/time.h (+4-10)
  • (modified) libc/cmake/modules/LLVMLibCHeaderRules.cmake (-15)
  • (modified) libc/docs/dev/header_generation.rst (+1-1)
  • (removed) libc/utils/hdrgen/hdrgen/gpu_headers.py (-54)
  • (modified) libc/utils/hdrgen/hdrgen/yaml_to_classes.py (+2-8)
diff --git a/clang/lib/Headers/llvm_libc_wrappers/assert.h b/clang/lib/Headers/llvm_libc_wrappers/assert.h
index 610ed96a458c6..7eadb2c354aab 100644
--- a/clang/lib/Headers/llvm_libc_wrappers/assert.h
+++ b/clang/lib/Headers/llvm_libc_wrappers/assert.h
@@ -19,13 +19,11 @@
 
 #if defined(__HIP__) || defined(__CUDA__)
 #define __LIBC_ATTRS __attribute__((device))
+#else
+#define __LIBC_ATTRS
 #endif
 
-#pragma omp begin declare target
-
-#include <llvm-libc-decls/assert.h>
-
-#pragma omp end declare target
+// TODO: Define these for CUDA / HIP.
 
 #undef __LIBC_ATTRS
 
diff --git a/clang/lib/Headers/llvm_libc_wrappers/ctype.h b/clang/lib/Headers/llvm_libc_wrappers/ctype.h
index 960cf43302c4c..79b0c1e9be953 100644
--- a/clang/lib/Headers/llvm_libc_wrappers/ctype.h
+++ b/clang/lib/Headers/llvm_libc_wrappers/ctype.h
@@ -13,128 +13,16 @@
 #error "This file is for GPU offloading compilation only"
 #endif
 
-// The GNU headers like to define 'toupper' and 'tolower' redundantly. This is
-// necessary to prevent it from doing that and remapping our implementation.
-#if (defined(__NVPTX__) || defined(__AMDGPU__)) && defined(__GLIBC__)
-#pragma push_macro("__USE_EXTERN_INLINES")
-#undef __USE_EXTERN_INLINES
-#endif
-
 #include_next <ctype.h>
 
-#if (defined(__NVPTX__) || defined(__AMDGPU__)) && defined(__GLIBC__)
-#pragma pop_macro("__USE_EXTERN_INLINES")
-#endif
-
-#if __has_include(<llvm-libc-decls/ctype.h>)
-
 #if defined(__HIP__) || defined(__CUDA__)
 #define __LIBC_ATTRS __attribute__((device))
+#else
+#define __LIBC_ATTRS
 #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("isascii")
-#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("toascii")
-#pragma push_macro("tolower")
-#pragma push_macro("toupper")
-#pragma push_macro("isalnum_l")
-#pragma push_macro("isalpha_l")
-#pragma push_macro("isascii_l")
-#pragma push_macro("isblank_l")
-#pragma push_macro("iscntrl_l")
-#pragma push_macro("isdigit_l")
-#pragma push_macro("isgraph_l")
-#pragma push_macro("islower_l")
-#pragma push_macro("isprint_l")
-#pragma push_macro("ispunct_l")
-#pragma push_macro("isspace_l")
-#pragma push_macro("isupper_l")
-#pragma push_macro("isxdigit_l")
-
-#undef isalnum
-#undef isalpha
-#undef isascii
-#undef iscntrl
-#undef isdigit
-#undef islower
-#undef isgraph
-#undef isprint
-#undef ispunct
-#undef isspace
-#undef isupper
-#undef isblank
-#undef isxdigit
-#undef toascii
-#undef tolower
-#undef toupper
-#undef isalnum_l
-#undef isalpha_l
-#undef iscntrl_l
-#undef isdigit_l
-#undef islower_l
-#undef isgraph_l
-#undef isprint_l
-#undef ispunct_l
-#undef isspace_l
-#undef isupper_l
-#undef isblank_l
-#undef isxdigit_l
-
-#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("isascii")
-#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("toascii")
-#pragma pop_macro("tolower")
-#pragma pop_macro("toupper")
-#pragma pop_macro("isalnum_l")
-#pragma pop_macro("isalpha_l")
-#pragma pop_macro("isascii_l")
-#pragma pop_macro("isblank_l")
-#pragma pop_macro("iscntrl_l")
-#pragma pop_macro("isdigit_l")
-#pragma pop_macro("isgraph_l")
-#pragma pop_macro("islower_l")
-#pragma pop_macro("isprint_l")
-#pragma pop_macro("ispunct_l")
-#pragma pop_macro("isspace_l")
-#pragma pop_macro("isupper_l")
-#pragma pop_macro("isxdigit_l")
-#endif
+// TODO: Define these for CUDA / HIP.
 
 #undef __LIBC_ATTRS
 
-#endif
-
 #endif // __CLANG_LLVM_LIBC_WRAPPERS_CTYPE_H__
diff --git a/clang/lib/Headers/llvm_libc_wrappers/inttypes.h b/clang/lib/Headers/llvm_libc_wrappers/inttypes.h
index 415f1e4b7bcab..22613898248f3 100644
--- a/clang/lib/Headers/llvm_libc_wrappers/inttypes.h
+++ b/clang/lib/Headers/llvm_libc_wrappers/inttypes.h
@@ -19,13 +19,11 @@
 
 #if defined(__HIP__) || defined(__CUDA__)
 #define __LIBC_ATTRS __attribute__((device))
+#else
+#define __LIBC_ATTRS
 #endif
 
-#pragma omp begin declare target
-
-#include <llvm-libc-decls/inttypes.h>
-
-#pragma omp end declare target
+// TODO: Define these for CUDA / HIP.
 
 #undef __LIBC_ATTRS
 
diff --git a/clang/lib/Headers/llvm_libc_wrappers/llvm-libc-decls/README.txt b/clang/lib/Headers/llvm_libc_wrappers/llvm-libc-decls/README.txt
deleted file mode 100644
index e012cd9e29312..0000000000000
--- a/clang/lib/Headers/llvm_libc_wrappers/llvm-libc-decls/README.txt
+++ /dev/null
@@ -1,6 +0,0 @@
-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.
diff --git a/clang/lib/Headers/llvm_libc_wrappers/stdio.h b/clang/lib/Headers/llvm_libc_wrappers/stdio.h
index 950f91b3763e8..0c3e44823da70 100644
--- a/clang/lib/Headers/llvm_libc_wrappers/stdio.h
+++ b/clang/lib/Headers/llvm_libc_wrappers/stdio.h
@@ -6,45 +6,19 @@
 //
 //===----------------------------------------------------------------------===//
 
+#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>
 
-// In some old versions of glibc, other standard headers sometimes define
-// special macros (e.g., __need_FILE) before including stdio.h to cause stdio.h
-// to produce special definitions.  Future includes of stdio.h when those
-// special macros are undefined are expected to produce the normal definitions
-// from stdio.h.
-//
-// We do not apply our include guard (__CLANG_LLVM_LIBC_WRAPPERS_STDIO_H__)
-// unconditionally to the above include_next.  Otherwise, after an occurrence of
-// the first glibc stdio.h use case described above, the include_next would be
-// skipped for remaining includes of stdio.h, leaving required symbols
-// undefined.
-//
-// We make the following assumptions to handle all use cases:
-//
-// 1. If the above include_next produces special glibc definitions, then (a) it
-//    does not produce the normal definitions that we must intercept below, (b)
-//    the current file was included from a glibc header that already defined
-//    __GLIBC__ (usually by including glibc's <features.h>), and (c) the above
-//    include_next does not define _STDIO_H.  In that case, we skip the rest of
-//    the current file and don't guard against future includes.
-// 2. If the above include_next produces the normal stdio.h definitions, then
-//    either (a) __GLIBC__ is not defined because C headers are from some other
-//    libc implementation or (b) the above include_next defines _STDIO_H to
-//    prevent the above include_next from having any effect in the future.
-#if !defined(__GLIBC__) || defined(_STDIO_H)
-
-#ifndef __CLANG_LLVM_LIBC_WRAPPERS_STDIO_H__
-#define __CLANG_LLVM_LIBC_WRAPPERS_STDIO_H__
-
-#if __has_include(<llvm-libc-decls/stdio.h>)
-
 #if defined(__HIP__) || defined(__CUDA__)
 #define __LIBC_ATTRS __attribute__((device))
+#else
+#define __LIBC_ATTRS
 #endif
 
 // Some headers provide these as macros. Temporarily undefine them so they do
@@ -60,21 +34,19 @@
 
 #pragma omp begin declare target
 
-#include <llvm-libc-decls/stdio.h>
+__LIBC_ATTRS extern FILE *stderr;
+__LIBC_ATTRS extern FILE *stdin;
+__LIBC_ATTRS extern FILE *stdout;
 
 #pragma omp end declare target
 
-#undef __LIBC_ATTRS
-
 // Restore the original macros when compiling on the host.
 #if !defined(__NVPTX__) && !defined(__AMDGPU__)
-#pragma pop_macro("stdout")
 #pragma pop_macro("stderr")
 #pragma pop_macro("stdin")
+#pragma pop_macro("stdout")
 #endif
 
-#endif
+#undef __LIBC_ATTRS
 
 #endif // __CLANG_LLVM_LIBC_WRAPPERS_STDIO_H__
-
-#endif
diff --git a/clang/lib/Headers/llvm_libc_wrappers/stdlib.h b/clang/lib/Headers/llvm_libc_wrappers/stdlib.h
index d79e7fa041ad4..7af5e2ebe031a 100644
--- a/clang/lib/Headers/llvm_libc_wrappers/stdlib.h
+++ b/clang/lib/Headers/llvm_libc_wrappers/stdlib.h
@@ -15,39 +15,18 @@
 
 #include_next <stdlib.h>
 
-#if __has_include(<llvm-libc-decls/stdlib.h>)
-
 #if defined(__HIP__) || defined(__CUDA__)
 #define __LIBC_ATTRS __attribute__((device))
+#else
+#define __LIBC_ATTRS
 #endif
 
 #pragma omp begin declare target
 
-// The LLVM C library uses these named types so we forward declare them.
-typedef void (*__atexithandler_t)(void);
-typedef int (*__search_compare_t)(const void *, const void *);
-typedef int (*__qsortcompare_t)(const void *, const void *);
-typedef int (*__qsortrcompare_t)(const void *, const void *, 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!");
-
-#if defined(__GLIBC__) && __cplusplus >= 201103L
-#define at_quick_exit atexit
-#endif
-
-#include <llvm-libc-decls/stdlib.h>
-
-#if defined(__GLIBC__) && __cplusplus >= 201103L
-#undef at_quick_exit
-#endif
+// TODO: Define these for CUDA / HIP.
 
 #pragma omp end declare target
 
 #undef __LIBC_ATTRS
 
-#endif
-
 #endif // __CLANG_LLVM_LIBC_WRAPPERS_STDLIB_H__
diff --git a/clang/lib/Headers/llvm_libc_wrappers/string.h b/clang/lib/Headers/llvm_libc_wrappers/string.h
index 0ea49cb137606..766a58f5b6db4 100644
--- a/clang/lib/Headers/llvm_libc_wrappers/string.h
+++ b/clang/lib/Headers/llvm_libc_wrappers/string.h
@@ -15,82 +15,14 @@
 
 #include_next <string.h>
 
-#if __has_include(<llvm-libc-decls/string.h>)
-
 #if defined(__HIP__) || defined(__CUDA__)
 #define __LIBC_ATTRS __attribute__((device))
-#endif
-
-#pragma omp begin declare target
-
-// The GNU headers provide C++ standard compliant headers when in C++ mode and
-// the LLVM libc does not. We need to manually provide the definitions using the
-// same prototypes.
-#if defined(__cplusplus) && defined(__GLIBC__) &&                              \
-    defined(__CORRECT_ISO_CPP_STRING_H_PROTO)
-
-#ifndef __LIBC_ATTRS
-#define __LIBC_ATTRS
-#endif
-
-extern "C" {
-void *memccpy(void *__restrict, const void *__restrict, int,
-              size_t) __LIBC_ATTRS;
-int memcmp(const void *, const void *, size_t) __LIBC_ATTRS;
-void *memcpy(void *__restrict, const void *__restrict, size_t) __LIBC_ATTRS;
-void *memmem(const void *, size_t, const void *, size_t) __LIBC_ATTRS;
-void *memmove(void *, const void *, size_t) __LIBC_ATTRS;
-void *mempcpy(void *__restrict, const void *__restrict, size_t) __LIBC_ATTRS;
-void *memset(void *, int, size_t) __LIBC_ATTRS;
-char *stpcpy(char *__restrict, const char *__restrict) __LIBC_ATTRS;
-char *stpncpy(char *__restrict, const char *__restrict, size_t) __LIBC_ATTRS;
-char *strcat(char *__restrict, const char *__restrict) __LIBC_ATTRS;
-int strcmp(const char *, const char *) __LIBC_ATTRS;
-int strcoll(const char *, const char *) __LIBC_ATTRS;
-char *strcpy(char *__restrict, const char *__restrict) __LIBC_ATTRS;
-size_t strcspn(const char *, const char *) __LIBC_ATTRS;
-char *strdup(const char *) __LIBC_ATTRS;
-size_t strlen(const char *) __LIBC_ATTRS;
-char *strncat(char *__restrict, const char *__restrict, size_t) __LIBC_ATTRS;
-int strncmp(const char *, const char *, size_t) __LIBC_ATTRS;
-char *strncpy(char *__restrict, const char *__restrict, size_t) __LIBC_ATTRS;
-char *strndup(const char *, size_t) __LIBC_ATTRS;
-size_t strnlen(const char *, size_t) __LIBC_ATTRS;
-size_t strspn(const char *, const char *) __LIBC_ATTRS;
-char *strtok(char *__restrict, const char *__restrict) __LIBC_ATTRS;
-char *strtok_r(char *__restrict, const char *__restrict,
-               char **__restrict) __LIBC_ATTRS;
-size_t strxfrm(char *__restrict, const char *__restrict, size_t) __LIBC_ATTRS;
-}
-
-extern "C++" {
-char *strstr(char *, const char *) noexcept __LIBC_ATTRS;
-const char *strstr(const char *, const char *) noexcept __LIBC_ATTRS;
-char *strpbrk(char *, const char *) noexcept __LIBC_ATTRS;
-const char *strpbrk(const char *, const char *) noexcept __LIBC_ATTRS;
-char *strrchr(char *, int) noexcept __LIBC_ATTRS;
-const char *strrchr(const char *, int) noexcept __LIBC_ATTRS;
-char *strchr(char *, int) noexcept __LIBC_ATTRS;
-const char *strchr(const char *, int) noexcept __LIBC_ATTRS;
-char *strchrnul(char *, int) noexcept __LIBC_ATTRS;
-const char *strchrnul(const char *, int) noexcept __LIBC_ATTRS;
-char *strcasestr(char *, const char *) noexcept __LIBC_ATTRS;
-const char *strcasestr(const char *, const char *) noexcept __LIBC_ATTRS;
-void *memrchr(void *__s, int __c, size_t __n) noexcept __LIBC_ATTRS;
-const void *memrchr(const void *__s, int __c, size_t __n) noexcept __LIBC_ATTRS;
-void *memchr(void *__s, int __c, size_t __n) noexcept __LIBC_ATTRS;
-const void *memchr(const void *__s, int __c, size_t __n) noexcept __LIBC_ATTRS;
-}
-
 #else
-#include <llvm-libc-decls/string.h>
-
+#define __LIBC_ATTRS
 #endif
 
-#pragma omp end declare target
+// TODO: Define these for CUDA / HIP.
 
 #undef __LIBC_ATTRS
 
-#endif
-
 #endif // __CLANG_LLVM_LIBC_WRAPPERS_STRING_H__
diff --git a/clang/lib/Headers/llvm_libc_wrappers/time.h b/clang/lib/Headers/llvm_libc_wrappers/time.h
index 9d1340c4eb748..d38eea327a199 100644
--- a/clang/lib/Headers/llvm_libc_wrappers/time.h
+++ b/clang/lib/Headers/llvm_libc_wrappers/time.h
@@ -15,20 +15,14 @@
 
 #include_next <time.h>
 
-#if __has_include(<llvm-libc-decls/time.h>)
-
 #if defined(__HIP__) || defined(__CUDA__)
 #define __LIBC_ATTRS __attribute__((device))
+#else
+#define __LIBC_ATTRS
 #endif
 
-#pragma omp begin declare target
-
-_Static_assert(sizeof(clock_t) == sizeof(long), "ABI mismatch!");
-
-#include <llvm-libc-decls/time.h>
+// TODO: Define these for CUDA / HIP.
 
-#pragma omp end declare target
-
-#endif
+#undef __LIBC_ATTRS
 
 #endif // __CLANG_LLVM_LIBC_WRAPPERS_TIME_H__
diff --git a/libc/cmake/modules/LLVMLibCHeaderRules.cmake b/libc/cmake/modules/LLVMLibCHeaderRules.cmake
index 01c288f0b9198..b092b37ffd8da 100644
--- a/libc/cmake/modules/LLVMLibCHeaderRules.cmake
+++ b/libc/cmake/modules/LLVMLibCHeaderRules.cmake
@@ -119,21 +119,6 @@ function(add_gen_header target_name)
     DEPFILE ${dep_file}
     COMMENT "Generating header ${ADD_GEN_HDR_GEN_HDR} from ${yaml_file}"
   )
-  if(LIBC_TARGET_OS_IS_GPU)
-    file(MAKE_DIRECTORY ${LIBC_INCLUDE_DIR}/llvm-libc-decls)
-    file(MAKE_DIRECTORY ${LIBC_INCLUDE_DIR}/llvm-libc-decls/gpu)
-    set(decl_out_file ${LIBC_INCLUDE_DIR}/llvm-libc-decls/${relative_path})
-    add_custom_command(
-      OUTPUT ${decl_out_file}
-      COMMAND ${Python3_EXECUTABLE} "${LIBC_SOURCE_DIR}/utils/hdrgen/yaml_to_classes.py"
-              ${yaml_file}
-              --export-decls
-              ${entry_points}
-              --output_dir ${decl_out_file}
-      WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}
-      DEPENDS ${yaml_file}
-    )
-  endif()
 
   if(ADD_GEN_HDR_DEPENDS)
     get_fq_deps_list(fq_deps_list ${ADD_GEN_HDR_DEPENDS})
diff --git a/libc/docs/dev/header_generation.rst b/libc/docs/dev/header_generation.rst
index 761a2822d6ab6..4cbad00e8ac98 100644
--- a/libc/docs/dev/header_generation.rst
+++ b/libc/docs/dev/header_generation.rst
@@ -131,7 +131,7 @@ Common Errors
 
      usage: yaml_to_classes.py [-h] [--output_dir OUTPUT_DIR] [--h_def_file H_DEF_FILE]
      [--add_function RETURN_TYPE NAME ARGUMENTS STANDARDS GUARD ATTRIBUTES]
-     [--e ENTRY_POINTS] [--export-decls]
+     [--e ENTRY_POINTS]
      yaml_file
      yaml_to_classes.py:
      error: argument --add_function: expected 6 arguments
diff --git a/libc/utils/hdrgen/hdrgen/gpu_headers.py b/libc/utils/hdrgen/hdrgen/gpu_headers.py
deleted file mode 100644
index 290bfe0649135..0000000000000
--- a/libc/utils/hdrgen/hdrgen/gpu_headers.py
+++ /dev/null
@@ -1,54 +0,0 @@
-# ===- GPU HeaderFile Class for --export-decls version --------*- python -*--==#
-#
-# 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
-#
-# ==-------------------------------------------------------------------------==#
-
-from hdrgen.header import HeaderFile
-
-
-class GpuHeaderFile(HeaderFile):
-    def __str__(self):
-        content = []
-
-        content.append(
-            f"//===-- C standard declarations for {self.name} ------------------------------===//"
-        )
-        content.append("//")
-        content.append(
-            "// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions."
-        )
-        content.append("// See https://llvm.org/LICENSE.txt for license information.")
-        content.append("// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception")
-        content.append("//")
-        content.append(
-            "//===----------------------------------------------------------------------===//\n"
-        )
-
-        header_guard = f"__LLVM_LIBC_DECLARATIONS_{self.name.upper()[:-2]}_H"
-        content.append(f"#ifndef {header_guard}")
-        content.append(f"#define {header_guard}\n")
-
-        content.append("#ifndef __LIBC_ATTRS")
-        content.append("#define __LIBC_ATTRS")
-        content.append("#endif\n")
-
-        content.append("#ifdef __cplusplus")
-        content.append('extern "C" {')
-        content.append("#endif\n")
-
-        for function in self.functions:
-            content.append(f"{function} __LIBC_ATTRS;\n")
-
-        for object in self.objects:
-            content.append(f"{object} __LIBC_ATTRS;\n")
-
-        content.append("#ifdef __cplusplus")
-        content.append("}")
-        content.append("#endif\n")
-
-        content.append(f"#endif")
-
-        return "\n".join(content)
diff --git a/libc/utils/hdrgen/hdrgen/yaml_to_classes.py b/libc/utils/hdrgen/hdrgen/yaml_to_classes.py
index 9eddbe615cbba..85aa3267b5274 100644
--- a/libc/utils/hdrgen/hdrgen/yaml_to_classes.py
+++ b/libc/utils/hdrgen/hdrgen/yaml_to_classes.py
@@ -14,7 +14,6 @@
 
 from hdrgen.enumeration import Enumeration
 from hdrgen.function import Function
-from hdrgen.gpu_headers import GpuHeaderFile as GpuHeader
 from hdrgen.header import HeaderFile
 from hdrgen.macro import Macro
 from hdrgen.object import Object
@@ -123,7 +122,7 @@ def load_yaml_file(yaml_file, header_class, entry_points):
 
     Args:
         yaml_file: Path to the YAML file.
-        header_class: The class to use for creating the header (HeaderFile or GpuHeader).
+        header_class: The class to use for creating the HeaderFile.
         entry_points: A list of specific function names to include in the header.
 
     Returns:
@@ -256,17 +255,12 @@ def main():
         help="Entry point to include",
         dest="entry_points",
     )
-    parser.add_argument(
-        "--export-decls",
-        action="store_true",
-        help="Flag to use GpuHeader for exporting declarations",
-    )
     args = parser.parse_args()
 
     if args.add_function:
         add_function_to_yaml(args.yaml_file, args.add_function)
 
-    header_class = GpuHeader if args.export_decls else HeaderFile
+    header_class = HeaderFile
     header = load_yaml_file(Path(args.yaml_file), header_class, args.entry_points)
 
     header_str = str(header)

Copy link
Contributor

@michaelrj-google michaelrj-google left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Libc changes LGTM

@github-actions
Copy link

🐧 Linux x64 Test Results

  • 111272 tests passed
  • 4418 tests skipped

@jhuber6
Copy link
Contributor Author

jhuber6 commented Nov 17, 2025

I could alternatively keep the decl handling for CUDA / HIP and just not do the include for OpenMP, hm. Somewhat curious if it would be worthwhile given SYCL. @sarnex any clue what declaring a device version of a function looks like in SYCL?

@sarnex
Copy link
Member

sarnex commented Nov 17, 2025

Not sure if I totally understand the question but if it's just how do we declare SYCL device functions basically it just comes down to slapping an attribute on the decl, I believe our plan for the upstream frontend is for it to be [[clang::sycl_external]]
If the code shouldn't be compiled on the host also it can be wrapped in a __SYCL_DEVICE_ONLY__ guard.

So if we want to compile externally callable device-only code for SYCL it would be something like

#ifdef __SYCL_DEVICE_ONLY__
[[clang::sycl_external]] void Foo();
#endif

@jhuber6
Copy link
Contributor Author

jhuber6 commented Nov 17, 2025

Not sure if I totally understand the question but if it's just how do we declare SYCL device functions basically it just comes down to slapping an attribute on the decl, I believe our plan for the upstream frontend is for it to be [[clang::sycl_external]] If the code shouldn't be compiled on the host also it can be wrapped in a __SYCL_DEVICE_ONLY__ guard.

So if we want to compile externally callable device-only code for SYCL it would be something like

#ifdef __SYCL_DEVICE_ONLY__
[[clang::sycl_external]] void Foo();
#endif

Yeah, the question is how to declare a function / global on the device side. In CUDA this is just the __device__ keyword / attribute. So the only way to do it is with macros?

@sarnex
Copy link
Member

sarnex commented Nov 17, 2025

Last I checked if you use just the attribute on the host the attr is dropped and the function is compiled as normal for the host, which we probably don't want.

@Fznamznon @tahonermann Do we have any plans to for a non-macro way to declare device-only code?

@Fznamznon
Copy link
Contributor

Fznamznon commented Nov 18, 2025

@Fznamznon @tahonermann Do we have any plans to for a non-macro way to declare device-only code?

That would be an extension to SYCL, I think.

In SYCL marking device code is not that obvious and mostly implicit. A function becomes device code if it is called from a kernel invoking API, i.e. parallel_for library call. We will use attribute [[clang::sycl_kernel_entry_point]] for marking these APIs in the SYCL library. A function marked with [[clang::sycl_kernel_entry_point]] will accept a callable object or a lambda and everything called from it recursively becomes device code. Additionally, there is SYCL_EXTERNAL macro which we will define via [[clang::sycl_external]] attribute. It simply forces code generation of a function for the device, even when it is not called from a callable object passed to [[clang::sycl_kernel_entry_point]] attributed function. The thing is, both these ways are forcing compilation for device meaning the host code will also have definitions for these functions. There is no way to define device-only function in SYCL except framing it with __SYCL_DEVICE_ONLY__ macro. It is the way SYCL 2020 is.

Copy link
Contributor

@DominikAdamski DominikAdamski left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM

@jhuber6 jhuber6 merged commit 7fe3564 into llvm:main Nov 19, 2025
32 checks passed
@llvm-ci
Copy link
Collaborator

llvm-ci commented Nov 19, 2025

LLVM Buildbot has detected a new failure on builder sanitizer-x86_64-linux-android running on sanitizer-buildbot-android while building clang,libc at step 2 "annotate".

Full details are available at: https://lab.llvm.org/buildbot/#/builders/186/builds/14095

Here is the relevant piece of the build log for the reference
Step 2 (annotate) failure: 'python ../sanitizer_buildbot/sanitizers/zorg/buildbot/builders/sanitizers/buildbot_selector.py' (failure)
...
[       OK ] AddressSanitizer.AtoiAndFriendsOOBTest (2272 ms)
[ RUN      ] AddressSanitizer.HasFeatureAddressSanitizerTest
[       OK ] AddressSanitizer.HasFeatureAddressSanitizerTest (0 ms)
[ RUN      ] AddressSanitizer.CallocReturnsZeroMem
[       OK ] AddressSanitizer.CallocReturnsZeroMem (16 ms)
[ DISABLED ] AddressSanitizer.DISABLED_TSDTest
[ RUN      ] AddressSanitizer.IgnoreTest
[       OK ] AddressSanitizer.IgnoreTest (0 ms)
[ RUN      ] AddressSanitizer.SignalTest
[       OK ] AddressSanitizer.SignalTest (211 ms)
[ RUN      ] AddressSanitizer.ReallocTest
[       OK ] AddressSanitizer.ReallocTest (34 ms)
[ RUN      ] AddressSanitizer.WrongFreeTest
[       OK ] AddressSanitizer.WrongFreeTest (117 ms)
[ RUN      ] AddressSanitizer.LongJmpTest
[       OK ] AddressSanitizer.LongJmpTest (0 ms)
[ RUN      ] AddressSanitizer.ThreadStackReuseTest
[       OK ] AddressSanitizer.ThreadStackReuseTest (7 ms)
[ DISABLED ] AddressSanitizer.DISABLED_MemIntrinsicUnalignedAccessTest
[ DISABLED ] AddressSanitizer.DISABLED_LargeFunctionSymbolizeTest
[ DISABLED ] AddressSanitizer.DISABLED_MallocFreeUnwindAndSymbolizeTest
[ RUN      ] AddressSanitizer.UseThenFreeThenUseTest
[       OK ] AddressSanitizer.UseThenFreeThenUseTest (140 ms)
[ RUN      ] AddressSanitizer.FileNameInGlobalReportTest
[       OK ] AddressSanitizer.FileNameInGlobalReportTest (105 ms)
[ DISABLED ] AddressSanitizer.DISABLED_StressStackReuseAndExceptionsTest
[ RUN      ] AddressSanitizer.MlockTest
[       OK ] AddressSanitizer.MlockTest (0 ms)
[ DISABLED ] AddressSanitizer.DISABLED_DemoThreadedTest
[ DISABLED ] AddressSanitizer.DISABLED_DemoStackTest
[ DISABLED ] AddressSanitizer.DISABLED_DemoThreadStackTest
[ DISABLED ] AddressSanitizer.DISABLED_DemoUAFLowIn
[ DISABLED ] AddressSanitizer.DISABLED_DemoUAFLowLeft
[ DISABLED ] AddressSanitizer.DISABLED_DemoUAFLowRight
[ DISABLED ] AddressSanitizer.DISABLED_DemoUAFHigh
[ DISABLED ] AddressSanitizer.DISABLED_DemoOOM
[ DISABLED ] AddressSanitizer.DISABLED_DemoDoubleFreeTest
[ DISABLED ] AddressSanitizer.DISABLED_DemoNullDerefTest
[ DISABLED ] AddressSanitizer.DISABLED_DemoFunctionStaticTest
[ DISABLED ] AddressSanitizer.DISABLED_DemoTooMuchMemoryTest
[ RUN      ] AddressSanitizer.LongDoubleNegativeTest
[       OK ] AddressSanitizer.LongDoubleNegativeTest (0 ms)
[----------] 19 tests from AddressSanitizer (27667 ms total)

[----------] Global test environment tear-down
[==========] 22 tests from 2 test suites ran. (27682 ms total)
[  PASSED  ] 22 tests.

  YOU HAVE 1 DISABLED TEST

Step 34 (run instrumented asan tests [aarch64/bluejay-userdebug/TQ3A.230805.001]) failure: run instrumented asan tests [aarch64/bluejay-userdebug/TQ3A.230805.001] (failure)
...
[ RUN      ] AddressSanitizer.HasFeatureAddressSanitizerTest
[       OK ] AddressSanitizer.HasFeatureAddressSanitizerTest (0 ms)
[ RUN      ] AddressSanitizer.CallocReturnsZeroMem
[       OK ] AddressSanitizer.CallocReturnsZeroMem (16 ms)
[ DISABLED ] AddressSanitizer.DISABLED_TSDTest
[ RUN      ] AddressSanitizer.IgnoreTest
[       OK ] AddressSanitizer.IgnoreTest (0 ms)
[ RUN      ] AddressSanitizer.SignalTest
[       OK ] AddressSanitizer.SignalTest (211 ms)
[ RUN      ] AddressSanitizer.ReallocTest
[       OK ] AddressSanitizer.ReallocTest (34 ms)
[ RUN      ] AddressSanitizer.WrongFreeTest
[       OK ] AddressSanitizer.WrongFreeTest (117 ms)
[ RUN      ] AddressSanitizer.LongJmpTest
[       OK ] AddressSanitizer.LongJmpTest (0 ms)
[ RUN      ] AddressSanitizer.ThreadStackReuseTest
[       OK ] AddressSanitizer.ThreadStackReuseTest (7 ms)
[ DISABLED ] AddressSanitizer.DISABLED_MemIntrinsicUnalignedAccessTest
[ DISABLED ] AddressSanitizer.DISABLED_LargeFunctionSymbolizeTest
[ DISABLED ] AddressSanitizer.DISABLED_MallocFreeUnwindAndSymbolizeTest
[ RUN      ] AddressSanitizer.UseThenFreeThenUseTest
[       OK ] AddressSanitizer.UseThenFreeThenUseTest (140 ms)
[ RUN      ] AddressSanitizer.FileNameInGlobalReportTest
[       OK ] AddressSanitizer.FileNameInGlobalReportTest (105 ms)
[ DISABLED ] AddressSanitizer.DISABLED_StressStackReuseAndExceptionsTest
[ RUN      ] AddressSanitizer.MlockTest
[       OK ] AddressSanitizer.MlockTest (0 ms)
[ DISABLED ] AddressSanitizer.DISABLED_DemoThreadedTest
[ DISABLED ] AddressSanitizer.DISABLED_DemoStackTest
[ DISABLED ] AddressSanitizer.DISABLED_DemoThreadStackTest
[ DISABLED ] AddressSanitizer.DISABLED_DemoUAFLowIn
[ DISABLED ] AddressSanitizer.DISABLED_DemoUAFLowLeft
[ DISABLED ] AddressSanitizer.DISABLED_DemoUAFLowRight
[ DISABLED ] AddressSanitizer.DISABLED_DemoUAFHigh
[ DISABLED ] AddressSanitizer.DISABLED_DemoOOM
[ DISABLED ] AddressSanitizer.DISABLED_DemoDoubleFreeTest
[ DISABLED ] AddressSanitizer.DISABLED_DemoNullDerefTest
[ DISABLED ] AddressSanitizer.DISABLED_DemoFunctionStaticTest
[ DISABLED ] AddressSanitizer.DISABLED_DemoTooMuchMemoryTest
[ RUN      ] AddressSanitizer.LongDoubleNegativeTest
[       OK ] AddressSanitizer.LongDoubleNegativeTest (0 ms)
[----------] 19 tests from AddressSanitizer (27667 ms total)

[----------] Global test environment tear-down
[==========] 22 tests from 2 test suites ran. (27682 ms total)
[  PASSED  ] 22 tests.

  YOU HAVE 1 DISABLED TEST
program finished with exit code 0
elapsedTime=2339.942983

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

backend:X86 clang:headers Headers provided by Clang, e.g. for intrinsics clang Clang issues not falling into any other category libc

Projects

None yet

Development

Successfully merging this pull request may close these issues.

7 participants