He escrito un kernel increíble que me traerá fama y fortuna, si solo puedo compilarlo con NVRTC:

#include <stdio.h>

__global__ void do_stuff() { }

Esperaba que los compiladores (en tiempo de ejecución) reconocieran los encabezados del sistema, como un compilador normal, y que esto "simplemente funcionaría" (módulo de cualquier maquinaria específica de printf). Alternativamente, si no funcionara, hubiera esperado un mensaje de error sobre la fuente de stdio.h no disponible con la llamada API "creación de programa" (nvrtcCreateProgram()), ya que paso NULL y NULL como sus dos últimos argumentos

Sin embargo, lo que obtengo es lo siguiente:

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

Esto me parece extraño. Significa que el compilador de tiempo de ejecución es capaz de buscar dentro de los encabezados del sistema, pero no puede encontrar stddef.h, como nvcc o el compilador del lado del host pueden .

¿Por qué sucede esto y cuál es la solución idiomática / recomendada?

Nota: Me gustaría una solución alternativa que sería multiplataforma, no solo trabajar en mi máquina individual.

0
einpoklum 12 jul. 2020 a las 14:42

2 respuestas

La mejor respuesta

Se toma un enfoque adicional en la biblioteca "JITify" que Robert Crovella me ha recordado gentilmente. Si bien esto no parece ser un documento muy bueno, Jitify incluye previamente fragmentos procesados de varios encabezados que considera adecuados. En particular para <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
)";

Para 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";

Y para 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";

Si incluye esas cadenas como encabezados, con los nombres apropiados como claves, es probable que su núcleo compile.

De hecho, uno podría formar archivos de encabezado a partir de estos y otros mini encabezados en jitify.hpp, para usar en la compilación del núcleo no NVRTC. Eso también puede ser útil.

Un último punto: las constantes anteriores no especifican un espacio de ejecución __device__. Entonces, o agrega __device__ allí, o le dice al compilador que asuma que las funciones están diseñadas para ejecutarse en el dispositivo solamente, a menos que se especifique lo contrario; esa es la opción del compilador --device-as-default-execution-space NVRTC.

0
einpoklum 25 jul. 2020 a las 22:33

Aquí hay dos soluciones que podrían funcionar, pero que preferiría evitar . Si son el único curso de acción razonable después de todo, comente y dígalo:

  1. Agregue la ruta específica a stddef.h como parámetro del compilador (-I o --include-path=).
  2. Pase la fuente de stddef.h a la llamada nvrtcCreateProgram().
0
einpoklum 12 jul. 2020 a las 11:42