Как правильно использовать include stdio.h в программе, скомпилированной с помощью NVRTC?

Я написал потрясающее ядро, которое принесет мне славу и богатство - если только я смогу собрать его с помощью NVRTC:

#include <stdio.h>

__global__ void do_stuff() { }

Я бы надеялся, что системные заголовки должны распознаваться (исполняемым) компилятором, как и обычный компилятор, и что это будет "просто работать" (по модулю любого механизма, специфичного для printf). В качестве альтернативы, если бы это не сработало, я ожидал сообщения об ошибкеstdio.hисточник недоступен с помощью вызова API "создания программы" (nvrtcCreateProgram()), поскольку я прохожу NULL а также NULL в качестве последних двух аргументов.

Однако я получаю следующее:

/usr/include/stdio.h(33): catastrophic error: cannot open source file "stddef.h"

Мне это кажется странным. Это означает, что во время выполнения компилятор находится в состоянии смотреть внутри системных заголовков, но не смог найтиstddef.h, например nvcc или компилятор на стороне хоста.

Почему это происходит и каков идиоматический / рекомендуемый обходной путь?

Примечание: мне бы хотелось, чтобы обходной путь был кроссплатформенным, а не работал только на моей индивидуальной машине.

Источник

2 ответа

Решение

Дополнительный подход используется в библиотеке "JITify", о которой мне любезно напомнил Роберт Кровелла. Хотя это, похоже, не очень хорошо документировано, Jitify предварительно включает обработанные фрагменты различных заголовков, которые он считает подходящими. В частности для<climits>/<limits.h>:

static const char* jitsafe_header_limits_h = R"(
#pragma once
#if defined _WIN32 || defined _WIN64
 #define __WORDSIZE 32
#else
 #if defined __x86_64__ && !defined __ILP32__
  #define __WORDSIZE 64
 #else
  #define __WORDSIZE 32
 #endif
#endif
#define MB_LEN_MAX  16
#define CHAR_BIT    8
#define SCHAR_MIN   (-128)
#define SCHAR_MAX   127
#define UCHAR_MAX   255
enum {
  _JITIFY_CHAR_IS_UNSIGNED = (char)-1 >= 0,
  CHAR_MIN = _JITIFY_CHAR_IS_UNSIGNED ? 0 : SCHAR_MIN,
  CHAR_MAX = _JITIFY_CHAR_IS_UNSIGNED ? UCHAR_MAX : SCHAR_MAX,
};
#define SHRT_MIN    (-32768)
#define SHRT_MAX    32767
#define USHRT_MAX   65535
#define INT_MIN     (-INT_MAX - 1)
#define INT_MAX     2147483647
#define UINT_MAX    4294967295U
#if __WORDSIZE == 64
 # define LONG_MAX  9223372036854775807L
#else
 # define LONG_MAX  2147483647L
#endif
#define LONG_MIN    (-LONG_MAX - 1L)
#if __WORDSIZE == 64
 #define ULONG_MAX  18446744073709551615UL
#else
 #define ULONG_MAX  4294967295UL
#endif
#define LLONG_MAX  9223372036854775807LL
#define LLONG_MIN  (-LLONG_MAX - 1LL)
#define ULLONG_MAX 18446744073709551615ULL
)";

за stddef.h:

static const char* jitsafe_header_stddef_h =
    "#pragma once\n"
    "#include <climits>\n"
    "namespace __jitify_stddef_ns {\n"
    "#if __cplusplus >= 201103L\n"
    "typedef decltype(nullptr) nullptr_t;\n"
    "#if defined(_MSC_VER)\n"
    "  typedef double max_align_t;\n"
    "#elif defined(__APPLE__)\n"
    "  typedef long double max_align_t;\n"
    "#else\n"
    "  // Define max_align_t to match the GCC definition.\n"
    "  typedef struct {\n"
    "    long long __jitify_max_align_nonce1\n"
    "        __attribute__((__aligned__(__alignof__(long long))));\n"
    "    long double __jitify_max_align_nonce2\n"
    "        __attribute__((__aligned__(__alignof__(long double))));\n"
    "  } max_align_t;\n"
    "#endif\n"
    "#endif  // __cplusplus >= 201103L\n"
    "#if __cplusplus >= 201703L\n"
    "enum class byte : unsigned char {};\n"
    "#endif  // __cplusplus >= 201703L\n"
    "} // namespace __jitify_stddef_ns\n"
    "namespace std {\n"
    "  // NVRTC provides built-in definitions of ::size_t and ::ptrdiff_t.\n"
    "  using ::size_t;\n"
    "  using ::ptrdiff_t;\n"
    "  using namespace __jitify_stddef_ns;\n"
    "} // namespace std\n"
    "using namespace __jitify_stddef_ns;\n";

и для stdio.h:

static const char* jitsafe_header_stdio_h =
    "#pragma once\n"
    "#include <stddef.h>\n"
    "#define FILE int\n"
    "int fflush ( FILE * stream );\n"
    "int fprintf ( FILE * stream, const char * format, ... );\n";

Если вы включите эти строки в качестве заголовков с соответствующими именами в качестве ключей, вполне вероятно, что ваше ядро ​​будет компилироваться.

Фактически, из этих и других мини-заголовков можно было формировать файлы заголовков в jitify.hpp, для использования при компиляции ядра, отличного от NVRTC. Это тоже может быть полезно.

И последнее: константы выше не определяют __device__пространство для исполнения. Итак, либо вы добавляете__device__там, или скажите компилятору предположить, что функции предназначены для выполнения только на устройстве, если не указано иное; это--device-as-default-execution-space Параметр компилятора NVRTC.

Источник

Вот два решения, которые могут сработать, но я бы предпочел избегать их. Если это, в конце концов, единственный разумный образ действий - прокомментируйте и скажите об этом:

  1. Добавьте конкретный путь к stddef.h как параметр компилятора (-I или --include-path=).
  2. Передайте источник stddef.h к nvrtcCreateProgram() вызов.
Источник
Другие вопросы по тегам