如何在CUDA NVRTC运行时编译代码中引入C标准头文件？

我之前在使用CUDA 9.2搭配NVRTC 7.5时也碰到过一模一样的问题，本质是老版本NVRTC对未标注执行空间的函数处理很严格：它默认把所有没加__host__/__device__/__global__的函数当成host函数，而运行时编译（JIT）模式下完全不允许包含host函数。当你引入<stdint.h>时，头文件里的一些辅助逻辑（比如类型检查宏、隐式的工具函数）没有被标记为__device__，直接触发了那个报错。

这里给你两个靠谱的解决方案，按需选择：

方案一：手动定义标准整数类型（最稳妥） #

既然问题根源是<stdint.h>的引入，我们可以直接绕开它，根据PTX的整数类型约定手动定义需要的类型。PTX的整数大小是固定的，和主机平台无关：

• 8位整数：对应char/unsigned char
• 16位整数：对应short/unsigned short
• 32位整数：对应int/unsigned int
• 64位整数：对应long long/unsigned long long

你只需要在核函数代码开头加上这些typedef：

// 手动替代<stdint.h>的核心类型定义，完全避开头文件问题
typedef signed char int8_t;
typedef unsigned char uint8_t;
typedef short int16_t;
typedef unsigned short uint16_t;
typedef int int32_t;
typedef unsigned int uint32_t;
typedef long long int64_t;
typedef unsigned long long uint64_t;

// 可选：如果需要用到类型极值宏，也可以手动定义
#define INT32_MIN (-2147483647 - 1)
#define INT32_MAX 2147483647
#define UINT32_MAX 0xFFFFFFFF

这样既可以正常使用int32_t等类型，又不会触发NVRTC的host函数报错，你的__global__函数也会正常编译成PTX里的.visible .entry f，完全不影响主机端调用。

方案二：强制让<stdint.h>内容适配设备端（适合必须用头文件的场景） #

如果你一定要引入<stdint.h>，可以通过临时重定义宏的方式，让头文件里的所有host函数自动转为device函数：

// 临时把__host__宏替换为__device__，让<stdint.h>里的函数都适配设备端
#define __host__ __device__
#include <stdint.h>
// 用完后立即还原__host__宏，避免影响后续代码
#undef __host__

extern "C" __global__ void f(int32_t* in, int32_t* out) {
    out[threadIdx.x] = in[threadIdx.x];
}

这个技巧会强制<stdint.h>里所有标注为__host__的函数都变成__device__，从根本上消除NVRTC对host函数的检测报错。而且不需要添加-default-device选项，你的__global__函数会保持.entry标记，正常从主机端调用。

不过要注意：这个方法可能有潜在副作用——如果<stdint.h>里某些函数原本只设计在主机端运行，强制转为device函数可能会出现其他编译问题，但在大多数使用标准整数类型的简单场景下，这个方法是完全可行的。

验证方法 #

• 方案一：把手动定义的typedef加到核函数源码里，去掉#include <stdint.h>，编译后查看PTX，会看到.visible .entry f的正确标记。
• 方案二：按上面的代码修改核函数源码，不需要调整编译选项，编译成功后同样能看到PTX里的.entry标记，且可以正常从主机端调用该核函数。

内容的提问来源于stack exchange，提问作者tmlen