我写了一个令人惊叹的内核,它将给我带来名声和财富-如果我只能用NVRTC编译它:
#include <stdio.h>
__global__ void do_stuff() { }
我希望系统头应该像常规编译器一样被(运行时(编译器识别,并且这将";只是工作;(模块化任何打印特定的机器(。或者,如果它不起作用,我会期望一条关于stdio.h
的源代码不可用的错误消息;程序创建";API调用(nvrtcCreateProgram()
(,因为我传递NULL
和NULL
作为它的最后两个参数。
然而,我得到的是:
/usr/include/stdio.h(33): catastrophic error: cannot open source file "stddef.h"
这对我来说似乎很奇怪。这意味着运行时编译器能够查看系统标头内部,但不能像nvcc或主机端编译器那样找到stddef.h
。
为什么会发生这种情况,惯用/推荐的解决方法是什么?
注意:我想要一个跨平台的解决方案,而不仅仅是在我的个人机器上工作。
;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 oncen"
"#include <climits>n"
"namespace __jitify_stddef_ns {n"
"#if __cplusplus >= 201103Ln"
"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"
"#elsen"
" // Define max_align_t to match the GCC definition.n"
" typedef struct {n"
" long long __jitify_max_align_nonce1n"
" __attribute__((__aligned__(__alignof__(long long))));n"
" long double __jitify_max_align_nonce2n"
" __attribute__((__aligned__(__alignof__(long double))));n"
" } max_align_t;n"
"#endifn"
"#endif // __cplusplus >= 201103Ln"
"#if __cplusplus >= 201703Ln"
"enum class byte : unsigned char {};n"
"#endif // __cplusplus >= 201703Ln"
"} // namespace __jitify_stddef_nsn"
"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 stdn"
"using namespace __jitify_stddef_ns;n";
对于stdio.h
:
static const char* jitsafe_header_stdio_h =
"#pragma oncen"
"#include <stddef.h>n"
"#define FILE intn"
"int fflush ( FILE * stream );n"
"int fprintf ( FILE * stream, const char * format, ... );n";
如果您将这些字符串作为头,并将适当的名称作为键,那么您的内核很可能会进行编译。
事实上,可以在jitify.hpp
中从这些和其他小型头文件中形成头文件,用于非NVRTC内核编译。这可能也很有用。
最后一点:上面的常量没有指定__device__
执行空间。因此,您可以在其中添加__device__
,或者告诉编译器假设函数仅用于在设备上执行,除非另有指定;这是--device-as-default-execution-space
NVRTC编译器选项。
这里有两个可能有效的解决方案,但我宁愿避免。如果它们毕竟是唯一合理的行动方案,请评论并这样说:
- 将
stddef.h
的特定路径添加为编译器参数(-I
或--include-path=
( - 将
stddef.h
的源传递给nvrtcCreateProgram()
调用