如何在 NVRTC 编译程序中正确使用 include stdio.h?
Posted
技术标签:
【中文标题】如何在 NVRTC 编译程序中正确使用 include stdio.h?【英文标题】:How to properly use include stdio.h in an NVRTC-compiled program? 【发布时间】:2020-11-01 18:00:06 【问题描述】:我编写了一个令人惊叹的内核,它将为我带来名利——如果我能用 NVRTC 编译它:
#include <stdio.h>
__global__ void do_stuff()
我希望系统头文件应该被(运行时)编译器识别,就像常规编译器一样,并且这将“正常工作”(以任何特定于 printf 的机器为模)。或者,如果它不起作用,我会期待一条关于 stdio.h
的源无法通过“程序创建”API 调用 (nvrtcCreateProgram()
) 提供的错误消息,因为我正在传递 NULL
和 @ 987654328@ 作为它的最后两个参数。
但是,我得到的是以下内容:
/usr/include/stdio.h(33): catastrophic error: cannot open source file "stddef.h"
这对我来说似乎很奇怪。这意味着运行时编译器能够查看系统头文件,但不能够找到stddef.h
,就像 nvcc 或主机端编译器一样。
为什么会发生这种情况,惯用/推荐的解决方法是什么?
注意:我想要一个跨平台的解决方法,而不仅仅是在我的个人机器上工作。
【问题讨论】:
this 可能感兴趣 也相关——***.com/q/50565200/681865——nvrtc 不应该真正与标准库头文件一起使用。使用 jitify,破解 stdlib 的功能集(不可移植),或者根本不使用它们 @talonmies:但你需要stdio.h
用于内核printf()
。尽管坦率地说,对我来说,包含的需要总是非常可疑,因为您实际上并没有使用标准库的 printf()
。
@RobertCrovella:jitify
是个好主意,但只是部分。你看,它有点像一个独立的“通天塔”;它为您简化了事情,不会与您可能正在使用的其他事情相结合;更不用说是一个 4K 行的头文件了。然而,在几个月内,我希望我能够提供更好的东西,或者至少是 NVIDIA 驱动程序 + nvrtc 基础架构,足以在上面编写类似 jitify 的功能,而且代码少得多。
我的意思是,nvrtc 真正适用于“纯”和“简单”的 CUDA 代码,无需 nvcc 前端处理的大量标准库解析和操作即可编译。我永远不会在我要 JIT 编译的东西中使用 printf
【参考方案1】:
这里有两种可能可行的解决方案,但我宁愿避免。如果它们毕竟是唯一合理的行动方案 - 请发表评论并这样说:
-
将特定路径添加到
stddef.h
作为编译器参数(-I
或--include-path=
)。
将stddef.h
的来源传递给nvrtcCreateProgram()
调用。
【讨论】:
【参考方案2】:“JITify”库中采用了另一种方法,Robert Crovella 亲切地提醒了我。虽然这似乎不是很好的文档,但 Jitify 预先包含了它认为适合的各种标头的已处理 sn-ps。特别是<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 编译器选项。
【讨论】:
以上是关于如何在 NVRTC 编译程序中正确使用 include stdio.h?的主要内容,如果未能解决你的问题,请参考以下文章