在 CUDA NVRTC 代码中包含 C 标准 headers
Including C standard headers in CUDA NVRTC code
我正在编写一个使用 NVRTC（CUDA 9.2 版和 NVRTC 7.5 版）在运行时编译的 CUDA 内核，它需要 stdint.h header，以便 int32_t等类型。
如果我编写没有包含的内核源代码，它可以正常工作。例如内核
extern "C" __global__ void f() { ... }
编译为 PTX 代码，其中 f 定义为 .visible .entry f。
但是如果内核源码是
#include <stdint.h>
extern "C" __global__ void f() { ... }
它报告 A function without execution space annotations (__host__/__device__/__global__) is considered a host function, and host functions are not allowed in JIT mode.（也没有 extern "C"）。
传递 -default-device 生成 PTX 代码 .visible .func f，因此无法从主机调用函数。
有没有办法在源代码中包含headers，并且仍然有__global__入口函数？或者，一种了解 NVRTC 编译器使用哪种整数大小约定的方法，以便可以手动定义 int32_t 等类型？
编辑：
显示问题的示例程序：
#include <cstdlib>
#include <string>
#include <vector>
#include <memory>
#include <cassert>
#include <iostream>
#include <cuda.h>
#include <cuda_runtime.h>
#include <nvrtc.h>
[[noreturn]] void fail(const std::string& msg, int code) {
std::cerr << "error: " << msg << " (" << code << ')' << std::endl;
std::exit(EXIT_FAILURE);
}
std::unique_ptr<char[]> compile_to_ptx(const char* program_source) {
nvrtcResult rv;
// create nvrtc program
nvrtcProgram prog;
rv = nvrtcCreateProgram(
&prog,
program_source,
"program.cu",
0,
nullptr,
nullptr
);
if(rv != NVRTC_SUCCESS) fail("nvrtcCreateProgram", rv);
// compile nvrtc program
std::vector<const char*> options = {
"--gpu-architecture=compute_30"
};
//options.push_back("-default-device");
rv = nvrtcCompileProgram(prog, options.size(), options.data());
if(rv != NVRTC_SUCCESS) {
std::size_t log_size;
rv = nvrtcGetProgramLogSize(prog, &log_size);
if(rv != NVRTC_SUCCESS) fail("nvrtcGetProgramLogSize", rv);
auto log = std::make_unique<char[]>(log_size);
rv = nvrtcGetProgramLog(prog, log.get());
if(rv != NVRTC_SUCCESS) fail("nvrtcGetProgramLog", rv);
assert(log[log_size - 1] == '[=13=]');
std::cerr << "Compile error; log:\n" << log.get() << std::endl;
fail("nvrtcCompileProgram", rv);
}
// get ptx code
std::size_t ptx_size;
rv = nvrtcGetPTXSize(prog, &ptx_size);
if(rv != NVRTC_SUCCESS) fail("nvrtcGetPTXSize", rv);
auto ptx = std::make_unique<char[]>(ptx_size);
rv = nvrtcGetPTX(prog, ptx.get());
if(rv != NVRTC_SUCCESS) fail("nvrtcGetPTX", rv);
assert(ptx[ptx_size - 1] == '[=13=]');
nvrtcDestroyProgram(&prog);
return ptx;
}
const char program_source[] = R"%%%(
//#include <stdint.h>
extern "C" __global__ void f(int* in, int* out) {
out[threadIdx.x] = in[threadIdx.x];
}
)%%%";
int main() {
CUresult rv;
// initialize CUDA
rv = cuInit(0);
if(rv != CUDA_SUCCESS) fail("cuInit", rv);
// compile program to ptx
auto ptx = compile_to_ptx(program_source);
std::cout << "PTX code:\n" << ptx.get() << std::endl;
}
当内核源代码中的 //#include <stdint.h> 取消注释时，它不再编译。当 //options.push_back("-default-device"); 取消注释时，它会编译但不会将函数 f 标记为 .entry.
CMakeLists.txt 编译它（需要 CUDA driver API + NVRTC）
cmake_minimum_required(VERSION 3.4)
project(cudabug CXX)
find_package(CUDA REQUIRED)
set(CMAKE_CXX_STANDARD 14)
set(CMAKE_CXX_STANDARD_REQUIRED 14)
add_executable(cudabug cudabug.cc)
include_directories(SYSTEM ${CUDA_INCLUDE_DIRS})
link_directories(${CUDA_LIBRARY_DIRS})
target_link_libraries(cudabug PUBLIC ${CUDA_LIBRARIES} nvrtc cuda)
[前言：这是一个非常hacky的答案，并且是GNU工具链特有的（虽然我怀疑问题中的问题也是GNU工具链特有的）]。
这里的问题似乎出在 GNU 标准头文件 features.h 中，它被拉入 stdint.h，然后最终定义了很多具有默认 [=14] 的存根函数=]编译space。这会导致 nvrtc 崩溃。 -default-device 选项似乎也会导致解析的 glibC 编译器功能集，这会使整个 nvrtc 编译器失败。
您可以通过为排除所有主机函数的标准库预定义一个功能集来解决这个问题（以一种非常 hacky 的方式）。将您的 JIT 内核代码更改为
const char program_source[] = R"%%%(
#define __ASSEMBLER__
#define __extension__
#include <stdint.h>
extern "C" __global__ void f(int32_t* in, int32_t* out) {
out[threadIdx.x] = in[threadIdx.x];
}
)%%%";
给我这个：
$ nvcc -std=c++14 -ccbin=g++-7 jit_header.cu -o jitheader -lnvrtc -lcuda
$ ./jitheader
PTX code:
//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-24330188
// Cuda compilation tools, release 9.2, V9.2.148
// Based on LLVM 3.4svn
//
.version 6.2
.target sm_30
.address_size 64
// .globl f
.visible .entry f(
.param .u64 f_param_0,
.param .u64 f_param_1
)
{
.reg .b32 %r<3>;
.reg .b64 %rd<8>;
ld.param.u64 %rd1, [f_param_0];
ld.param.u64 %rd2, [f_param_1];
cvta.to.global.u64 %rd3, %rd2;
cvta.to.global.u64 %rd4, %rd1;
mov.u32 %r1, %tid.x;
mul.wide.u32 %rd5, %r1, 4;
add.s64 %rd6, %rd4, %rd5;
ld.global.u32 %r2, [%rd6];
add.s64 %rd7, %rd3, %rd5;
st.global.u32 [%rd7], %r2;
ret;
}
重要警告：这在我试过的 glibC 系统上有效。它可能不适用于其他工具链或 libC 实现（如果确实存在此问题）。
另一种选择是为某些标准库 header 创建 stand-ins。 NVRTC 的 API 支持您将 header 文件内容指定为字符串，并与 header 名称相关联 - 在它为您查找文件系统之前。这种方法在 NVIDIA JITify 中被采用，我自己也采用了它来处理可能会或可能不会发布的其他内容。
执行此操作的简单方法您只需从 here 中获取 stdint.h、limits.h 的 JITify header 存根，我也附加了它，因为不是很长。或者，您可以自己生成此存根以确保您不会遗漏与标准相关的任何内容。工作原理如下：
从您的 stdint.h 文件开始（或 cstdint 文件，视情况而定）；
对于文件中的每个 include 指令（递归地，对于 include 等中的每个 include）：
2.1 确定是否可以完全跳过包含文件（可能通过进行一些已知在 GPU 上保留的定义）。
2.2 如果您不确定是否可以跳过该文件 - 将其完全包含并递归到 (2.)，或者将其单独保存 header（并在 (1. .) 到它）。
您现在有一个 header 文件，其中仅包含 device-safe header 个文件（或 none 个文件）
Partially-preprocess 文件，删除所有不会在 GPU 上使用的内容
删除在 GPU 上可能有问题的行（例如 #pragma's），并根据每个函数声明添加 __device__ __host__ 或 __host__ 。
重要说明：这样做需要注意许可证和版权。您将创建 glibc and/or JITify and/or Whosebug 贡献等的“衍生作品”
现在，我承诺的来自 NVIDIA JITify 的 stdint.h 和 limits.h。我已将它们修改为没有命名空间：
stdint.h:
#pragma once
#include <limits.h>
typedef signed char int8_t;
typedef signed short int16_t;
typedef signed int int32_t;
typedef signed long long int64_t;
typedef signed char int_fast8_t;
typedef signed short int_fast16_t;
typedef signed int int_fast32_t;
typedef signed long long int_fast64_t;
typedef signed char int_least8_t;
typedef signed short int_least16_t;
typedef signed int int_least32_t;
typedef signed long long int_least64_t;
typedef signed long long intmax_t;
typedef signed long intptr_t; //optional
typedef unsigned char uint8_t;
typedef unsigned short uint16_t;
typedef unsigned int uint32_t;
typedef unsigned long long uint64_t;
typedef unsigned char uint_fast8_t;
typedef unsigned short uint_fast16_t;
typedef unsigned int uint_fast32_t;
typedef unsigned long long uint_fast64_t;
typedef unsigned char uint_least8_t;
typedef unsigned short uint_least16_t;
typedef unsigned int uint_least32_t;
typedef unsigned long long uint_least64_t;
typedef unsigned long long uintmax_t;
#define INT8_MIN SCHAR_MIN
#define INT16_MIN SHRT_MIN
#if defined _WIN32 || defined _WIN64
#define WCHAR_MIN SHRT_MIN
#define WCHAR_MAX SHRT_MAX
typedef unsigned long long uintptr_t; //optional
#else
#define WCHAR_MIN INT_MIN
#define WCHAR_MAX INT_MAX
typedef unsigned long uintptr_t; //optional
#endif
#define INT32_MIN INT_MIN
#define INT64_MIN LLONG_MIN
#define INT8_MAX SCHAR_MAX
#define INT16_MAX SHRT_MAX
#define INT32_MAX INT_MAX
#define INT64_MAX LLONG_MAX
#define UINT8_MAX UCHAR_MAX
#define UINT16_MAX USHRT_MAX
#define UINT32_MAX UINT_MAX
#define UINT64_MAX ULLONG_MAX
#define INTPTR_MIN LONG_MIN
#define INTMAX_MIN LLONG_MIN
#define INTPTR_MAX LONG_MAX
#define INTMAX_MAX LLONG_MAX
#define UINTPTR_MAX ULONG_MAX
#define UINTMAX_MAX ULLONG_MAX
#define PTRDIFF_MIN INTPTR_MIN
#define PTRDIFF_MAX INTPTR_MAX
#define SIZE_MAX UINT64_MAX
limits.h:
#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
我正在编写一个使用 NVRTC（CUDA 9.2 版和 NVRTC 7.5 版）在运行时编译的 CUDA 内核，它需要 stdint.h header，以便 int32_t等类型。
如果我编写没有包含的内核源代码，它可以正常工作。例如内核
extern "C" __global__ void f() { ... }
编译为 PTX 代码，其中 f 定义为 .visible .entry f。
但是如果内核源码是
#include <stdint.h>
extern "C" __global__ void f() { ... }
它报告 A function without execution space annotations (__host__/__device__/__global__) is considered a host function, and host functions are not allowed in JIT mode.（也没有 extern "C"）。
传递 -default-device 生成 PTX 代码 .visible .func f，因此无法从主机调用函数。
有没有办法在源代码中包含headers，并且仍然有__global__入口函数？或者，一种了解 NVRTC 编译器使用哪种整数大小约定的方法，以便可以手动定义 int32_t 等类型？
编辑： 显示问题的示例程序：
#include <cstdlib>
#include <string>
#include <vector>
#include <memory>
#include <cassert>
#include <iostream>
#include <cuda.h>
#include <cuda_runtime.h>
#include <nvrtc.h>
[[noreturn]] void fail(const std::string& msg, int code) {
std::cerr << "error: " << msg << " (" << code << ')' << std::endl;
std::exit(EXIT_FAILURE);
}
std::unique_ptr<char[]> compile_to_ptx(const char* program_source) {
nvrtcResult rv;
// create nvrtc program
nvrtcProgram prog;
rv = nvrtcCreateProgram(
&prog,
program_source,
"program.cu",
0,
nullptr,
nullptr
);
if(rv != NVRTC_SUCCESS) fail("nvrtcCreateProgram", rv);
// compile nvrtc program
std::vector<const char*> options = {
"--gpu-architecture=compute_30"
};
//options.push_back("-default-device");
rv = nvrtcCompileProgram(prog, options.size(), options.data());
if(rv != NVRTC_SUCCESS) {
std::size_t log_size;
rv = nvrtcGetProgramLogSize(prog, &log_size);
if(rv != NVRTC_SUCCESS) fail("nvrtcGetProgramLogSize", rv);
auto log = std::make_unique<char[]>(log_size);
rv = nvrtcGetProgramLog(prog, log.get());
if(rv != NVRTC_SUCCESS) fail("nvrtcGetProgramLog", rv);
assert(log[log_size - 1] == '[=13=]');
std::cerr << "Compile error; log:\n" << log.get() << std::endl;
fail("nvrtcCompileProgram", rv);
}
// get ptx code
std::size_t ptx_size;
rv = nvrtcGetPTXSize(prog, &ptx_size);
if(rv != NVRTC_SUCCESS) fail("nvrtcGetPTXSize", rv);
auto ptx = std::make_unique<char[]>(ptx_size);
rv = nvrtcGetPTX(prog, ptx.get());
if(rv != NVRTC_SUCCESS) fail("nvrtcGetPTX", rv);
assert(ptx[ptx_size - 1] == '[=13=]');
nvrtcDestroyProgram(&prog);
return ptx;
}
const char program_source[] = R"%%%(
//#include <stdint.h>
extern "C" __global__ void f(int* in, int* out) {
out[threadIdx.x] = in[threadIdx.x];
}
)%%%";
int main() {
CUresult rv;
// initialize CUDA
rv = cuInit(0);
if(rv != CUDA_SUCCESS) fail("cuInit", rv);
// compile program to ptx
auto ptx = compile_to_ptx(program_source);
std::cout << "PTX code:\n" << ptx.get() << std::endl;
}
当内核源代码中的 //#include <stdint.h> 取消注释时，它不再编译。当 //options.push_back("-default-device"); 取消注释时，它会编译但不会将函数 f 标记为 .entry.
CMakeLists.txt 编译它（需要 CUDA driver API + NVRTC）
cmake_minimum_required(VERSION 3.4)
project(cudabug CXX)
find_package(CUDA REQUIRED)
set(CMAKE_CXX_STANDARD 14)
set(CMAKE_CXX_STANDARD_REQUIRED 14)
add_executable(cudabug cudabug.cc)
include_directories(SYSTEM ${CUDA_INCLUDE_DIRS})
link_directories(${CUDA_LIBRARY_DIRS})
target_link_libraries(cudabug PUBLIC ${CUDA_LIBRARIES} nvrtc cuda)
[前言：这是一个非常hacky的答案，并且是GNU工具链特有的（虽然我怀疑问题中的问题也是GNU工具链特有的）]。
这里的问题似乎出在 GNU 标准头文件 features.h 中，它被拉入 stdint.h，然后最终定义了很多具有默认 [=14] 的存根函数=]编译space。这会导致 nvrtc 崩溃。 -default-device 选项似乎也会导致解析的 glibC 编译器功能集，这会使整个 nvrtc 编译器失败。
您可以通过为排除所有主机函数的标准库预定义一个功能集来解决这个问题（以一种非常 hacky 的方式）。将您的 JIT 内核代码更改为
const char program_source[] = R"%%%(
#define __ASSEMBLER__
#define __extension__
#include <stdint.h>
extern "C" __global__ void f(int32_t* in, int32_t* out) {
out[threadIdx.x] = in[threadIdx.x];
}
)%%%";
给我这个：
$ nvcc -std=c++14 -ccbin=g++-7 jit_header.cu -o jitheader -lnvrtc -lcuda
$ ./jitheader
PTX code:
//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-24330188
// Cuda compilation tools, release 9.2, V9.2.148
// Based on LLVM 3.4svn
//
.version 6.2
.target sm_30
.address_size 64
// .globl f
.visible .entry f(
.param .u64 f_param_0,
.param .u64 f_param_1
)
{
.reg .b32 %r<3>;
.reg .b64 %rd<8>;
ld.param.u64 %rd1, [f_param_0];
ld.param.u64 %rd2, [f_param_1];
cvta.to.global.u64 %rd3, %rd2;
cvta.to.global.u64 %rd4, %rd1;
mov.u32 %r1, %tid.x;
mul.wide.u32 %rd5, %r1, 4;
add.s64 %rd6, %rd4, %rd5;
ld.global.u32 %r2, [%rd6];
add.s64 %rd7, %rd3, %rd5;
st.global.u32 [%rd7], %r2;
ret;
}
重要警告：这在我试过的 glibC 系统上有效。它可能不适用于其他工具链或 libC 实现（如果确实存在此问题）。
另一种选择是为某些标准库 header 创建 stand-ins。 NVRTC 的 API 支持您将 header 文件内容指定为字符串，并与 header 名称相关联 - 在它为您查找文件系统之前。这种方法在 NVIDIA JITify 中被采用，我自己也采用了它来处理可能会或可能不会发布的其他内容。
执行此操作的简单方法您只需从 here 中获取 stdint.h、limits.h 的 JITify header 存根，我也附加了它，因为不是很长。或者，您可以自己生成此存根以确保您不会遗漏与标准相关的任何内容。工作原理如下：
从您的
stdint.h文件开始（或cstdint文件，视情况而定）；对于文件中的每个 include 指令（递归地，对于 include 等中的每个 include）：
2.1 确定是否可以完全跳过包含文件（可能通过进行一些已知在 GPU 上保留的定义）。
2.2 如果您不确定是否可以跳过该文件 - 将其完全包含并递归到 (2.)，或者将其单独保存 header（并在 (1. .) 到它）。
您现在有一个 header 文件，其中仅包含 device-safe header 个文件（或 none 个文件）
Partially-preprocess 文件，删除所有不会在 GPU 上使用的内容 删除在 GPU 上可能有问题的行（例如
#pragma's），并根据每个函数声明添加__device____host__或__host__ 。
重要说明：这样做需要注意许可证和版权。您将创建 glibc and/or JITify and/or Whosebug 贡献等的“衍生作品”
现在，我承诺的来自 NVIDIA JITify 的 stdint.h 和 limits.h。我已将它们修改为没有命名空间：
stdint.h:
#pragma once
#include <limits.h>
typedef signed char int8_t;
typedef signed short int16_t;
typedef signed int int32_t;
typedef signed long long int64_t;
typedef signed char int_fast8_t;
typedef signed short int_fast16_t;
typedef signed int int_fast32_t;
typedef signed long long int_fast64_t;
typedef signed char int_least8_t;
typedef signed short int_least16_t;
typedef signed int int_least32_t;
typedef signed long long int_least64_t;
typedef signed long long intmax_t;
typedef signed long intptr_t; //optional
typedef unsigned char uint8_t;
typedef unsigned short uint16_t;
typedef unsigned int uint32_t;
typedef unsigned long long uint64_t;
typedef unsigned char uint_fast8_t;
typedef unsigned short uint_fast16_t;
typedef unsigned int uint_fast32_t;
typedef unsigned long long uint_fast64_t;
typedef unsigned char uint_least8_t;
typedef unsigned short uint_least16_t;
typedef unsigned int uint_least32_t;
typedef unsigned long long uint_least64_t;
typedef unsigned long long uintmax_t;
#define INT8_MIN SCHAR_MIN
#define INT16_MIN SHRT_MIN
#if defined _WIN32 || defined _WIN64
#define WCHAR_MIN SHRT_MIN
#define WCHAR_MAX SHRT_MAX
typedef unsigned long long uintptr_t; //optional
#else
#define WCHAR_MIN INT_MIN
#define WCHAR_MAX INT_MAX
typedef unsigned long uintptr_t; //optional
#endif
#define INT32_MIN INT_MIN
#define INT64_MIN LLONG_MIN
#define INT8_MAX SCHAR_MAX
#define INT16_MAX SHRT_MAX
#define INT32_MAX INT_MAX
#define INT64_MAX LLONG_MAX
#define UINT8_MAX UCHAR_MAX
#define UINT16_MAX USHRT_MAX
#define UINT32_MAX UINT_MAX
#define UINT64_MAX ULLONG_MAX
#define INTPTR_MIN LONG_MIN
#define INTMAX_MIN LLONG_MIN
#define INTPTR_MAX LONG_MAX
#define INTMAX_MAX LLONG_MAX
#define UINTPTR_MAX ULONG_MAX
#define UINTMAX_MAX ULLONG_MAX
#define PTRDIFF_MIN INTPTR_MIN
#define PTRDIFF_MAX INTPTR_MAX
#define SIZE_MAX UINT64_MAX
limits.h:
#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