在 CUDA NVRTC 代码中包含 C 标准头文件
Posted
技术标签:
【中文标题】在 CUDA NVRTC 代码中包含 C 标准头文件【英文标题】：Including C standard headers in CUDA NVRTC code 【发布时间】：2018-11-06 23:28:48 【问题描述】：我正在编写一个在运行时使用 NVRTC（CUDA 9.2 版和 NVRTC 7.5 版）编译的 CUDA 内核，它需要 stdint.h 标头，以便拥有 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，因此无法从主机调用该函数。
有没有办法在源代码中包含标头，并且仍然有__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] == '\0');
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] == '\0');
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 驱动 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)
【问题讨论】：
您是否有实际的复制案例可以编辑到您的问题中？ 如果只是关于 int32_t，使用 int32_t = int 应该可以解决问题。如果是关于标题的一般性问题，那么问题很广泛但非常有趣:) 添加示例程序 【参考方案1】：[前言：这是一个非常老套的答案，并且特定于 GNU 工具链（尽管我怀疑问题中的问题也特定于 GNU 工具链）]。
这里的问题似乎在于 GNU 标准头 features.h，它被拉入 stdint.h，然后最终定义了许多具有默认 __host__ 编译空间的存根函数。这会导致 nvrtc 炸毁。似乎-default-device 选项会导致已解析的 glibC 编译器功能集导致整个 nvrtc 编译器失败。
您可以通过为标准库预定义一个排除所有主机函数的功能集来解决这个问题（以一种非常老套的方式）。将 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 实现（如果他们确实有这个问题）。
【讨论】：
【参考方案2】：另一种选择是为某些标准库头文件创建替代项。 NVRTC 的 API 支持您将头文件内容指定为与头名称相关联的字符串 - 在它为您查看文件系统之前。 NVIDIA JITify 采用了这种方法，我自己也采用了这种方法来处理可能会或可能不会发布的其他内容。
执行此操作的简单方法您可以从here 获取stdint.h、limits.h 的JITify 标头存根，我也附上了它，因为它不是很长。或者，您可以自己生成此存根，以确保您不会错过与标准相关的任何内容。以下是它的工作原理：
从您的stdint.h 文件开始（或cstdint 文件，视情况而定）；
对于文件中的每个包含指令（递归地，对于包含等中的每个包含）：
2.1 确定是否可以完全跳过包含文件（可能通过进行一些已知在 GPU 上保存的定义）。
2.2 如果您不确定是否可以跳过该文件 - 将其完全包含并递归到 (2.)，或者将其保留为单独的标头（并将 (1.) 中的整个过程应用于它）。
你现在有一个头文件，它只包含设备安全的头文件（或根本没有）
对文件进行部分预处理，删除不会在 GPU 上使用的所有内容
删除在 GPU 上可能有问题的行（例如#pragma's），并根据每个函数声明添加__device__ __host__ 或仅__host__。
重要提示：这样做需要注意许可和版权。您将创建 glibc 和/或 JITify 和/或 *** 贡献等的“衍生作品”。
现在，我承诺了来自 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
【讨论】：
以上是关于在 CUDA NVRTC 代码中包含 C 标准头文件的主要内容，如果未能解决你的问题，请参考以下文章