c++ - NVCC 和 NVRTC 在编译为 PTX 时的区别

Question

总结

我正在移植一个基于 Scratchapixel version 的简单光线追踪应用程序到一堆 GPU 库。我使用运行时 API 和驱动程序 API 成功地将它移植到 CUDA，但它抛出一个 Segmentation fault (core dumped)当我尝试将在运行时编译的 PTX 与 NVRTC 一起使用时。如果我取消注释 #include <math.h>内核文件开头的指令(见下文)，它仍然可以使用 NVCC(生成的 PTX 完全相同)但在使用 NVRTC 编译时失败。

我想知道如何让 NVRTC 像 NVCC 一样运行(这可能吗？)，或者至少了解这个问题背后的原因。

详细说明

文件 kernel.cu (内核源代码):

//#include <math.h>

#define MAX_RAY_DEPTH 5

template<typename T>
class Vec3
{
public:
    T x, y, z;
    __device__ Vec3() : x(T(0)), y(T(0)), z(T(0)) {}
    __device__ Vec3(T xx) : x(xx), y(xx), z(xx) {}
    __device__ Vec3(T xx, T yy, T zz) : x(xx), y(yy), z(zz) {}
    __device__ Vec3& normalize()
    {
        T nor2 = length2();
        if (nor2 > 0) {
            T invNor = 1 / sqrt(nor2);
            x *= invNor, y *= invNor, z *= invNor;
        }
        return *this;
    }
    __device__ Vec3<T> operator * (const T &f) const { return Vec3<T>(x * f, y * f, z * f); }
    __device__ Vec3<T> operator * (const Vec3<T> &v) const { return Vec3<T>(x * v.x, y * v.y, z * v.z); }
    __device__ T dot(const Vec3<T> &v) const { return x * v.x + y * v.y + z * v.z; }
    __device__ Vec3<T> operator - (const Vec3<T> &v) const { return Vec3<T>(x - v.x, y - v.y, z - v.z); }
    __device__ Vec3<T> operator + (const Vec3<T> &v) const { return Vec3<T>(x + v.x, y + v.y, z + v.z); }
    __device__ Vec3<T>& operator += (const Vec3<T> &v) { x += v.x, y += v.y, z += v.z; return *this; }
    __device__ Vec3<T>& operator *= (const Vec3<T> &v) { x *= v.x, y *= v.y, z *= v.z; return *this; }
    __device__ Vec3<T> operator - () const { return Vec3<T>(-x, -y, -z); }
    __device__ T length2() const { return x * x + y * y + z * z; }
    __device__ T length() const { return sqrt(length2()); }
};

typedef Vec3<float> Vec3f;
typedef Vec3<bool> Vec3b;

class Sphere
{
public:
    const char* id;
    Vec3f center;                           /// position of the sphere
    float radius, radius2;                  /// sphere radius and radius^2
    Vec3f surfaceColor, emissionColor;      /// surface color and emission (light)
    float transparency, reflection;         /// surface transparency and reflectivity
    int animation_frame;
    Vec3b animation_position_rand;
    Vec3f animation_position;
    Sphere(
        const char* id,
        const Vec3f &c,
        const float &r,
        const Vec3f &sc,
        const float &refl = 0,
        const float &transp = 0,
        const Vec3f &ec = 0) :
        id(id), center(c), radius(r), radius2(r * r), surfaceColor(sc),
        emissionColor(ec), transparency(transp), reflection(refl)
    {
        animation_frame = 0;
    }
    //[comment]
    // Compute a ray-sphere intersection using the geometric solution
    //[/comment]
    __device__ bool intersect(const Vec3f &rayorig, const Vec3f &raydir, float &t0, float &t1) const
    {
        Vec3f l = center - rayorig;
        float tca = l.dot(raydir);
        if (tca < 0) return false;
        float d2 = l.dot(l) - tca * tca;
        if (d2 > radius2) return false;
        float thc = sqrt(radius2 - d2);
        t0 = tca - thc;
        t1 = tca + thc;

        return true;
    }
};

__device__ float mix(const float &a, const float &b, const float &mixval)
{
    return b * mixval + a * (1 - mixval);
}

__device__ Vec3f trace(
    const Vec3f &rayorig,
    const Vec3f &raydir,
    const Sphere *spheres,
    const unsigned int spheres_size,
    const int &depth)
{
    float tnear = INFINITY;
    const Sphere* sphere = NULL;
    // find intersection of this ray with the sphere in the scene
    for (unsigned i = 0; i < spheres_size; ++i) {
        float t0 = INFINITY, t1 = INFINITY;
        if (spheres[i].intersect(rayorig, raydir, t0, t1)) {
            if (t0 < 0) t0 = t1;
            if (t0 < tnear) {
                tnear = t0;
                sphere = &spheres[i];
            }
        }
    }
    // if there's no intersection return black or background color
    if (!sphere) return Vec3f(2);
    Vec3f surfaceColor = 0; // color of the ray/surfaceof the object intersected by the ray
    Vec3f phit = rayorig + raydir * tnear; // point of intersection
    Vec3f nhit = phit - sphere->center; // normal at the intersection point
    nhit.normalize(); // normalize normal direction
    // If the normal and the view direction are not opposite to each other
    // reverse the normal direction. That also means we are inside the sphere so set
    // the inside bool to true. Finally reverse the sign of IdotN which we want
    // positive.
    float bias = 1e-4; // add some bias to the point from which we will be tracing
    bool inside = false;
    if (raydir.dot(nhit) > 0) nhit = -nhit, inside = true;
    if ((sphere->transparency > 0 || sphere->reflection > 0) && depth < MAX_RAY_DEPTH) {
        float facingratio = -raydir.dot(nhit);
        // change the mix value to tweak the effect
        float fresneleffect = mix(pow(1 - facingratio, 3), 1, 0.1);
        // compute reflection direction (not need to normalize because all vectors
        // are already normalized)
        Vec3f refldir = raydir - nhit * 2 * raydir.dot(nhit);
        refldir.normalize();
        Vec3f reflection = trace(phit + nhit * bias, refldir, spheres, spheres_size, depth + 1);
        Vec3f refraction = 0;
        // if the sphere is also transparent compute refraction ray (transmission)
        if (sphere->transparency) {
            float ior = 1.1, eta = (inside) ? ior : 1 / ior; // are we inside or outside the surface?
            float cosi = -nhit.dot(raydir);
            float k = 1 - eta * eta * (1 - cosi * cosi);
            Vec3f refrdir = raydir * eta + nhit * (eta *  cosi - sqrt(k));
            refrdir.normalize();
            refraction = trace(phit - nhit * bias, refrdir, spheres, spheres_size, depth + 1);
        }
        // the result is a mix of reflection and refraction (if the sphere is transparent)
        surfaceColor = (
            reflection * fresneleffect +
            refraction * (1 - fresneleffect) * sphere->transparency) * sphere->surfaceColor;
    }
    else {
        // it's a diffuse object, no need to raytrace any further
        for (unsigned i = 0; i < spheres_size; ++i) {
            if (spheres[i].emissionColor.x > 0) {
                // this is a light
                Vec3f transmission = 1;
                Vec3f lightDirection = spheres[i].center - phit;
                lightDirection.normalize();
                for (unsigned j = 0; j < spheres_size; ++j) {
                    if (i != j) {
                        float t0, t1;
                        if (spheres[j].intersect(phit + nhit * bias, lightDirection, t0, t1)) {
                            transmission = 0;
                            break;
                        }
                    }
                }
                surfaceColor += sphere->surfaceColor * transmission *
                max(float(0), nhit.dot(lightDirection)) * spheres[i].emissionColor;
            }
        }
    }

    return surfaceColor + sphere->emissionColor;
}

extern "C" __global__
void raytrace_kernel(unsigned int width, unsigned int height, Vec3f *image, Sphere *spheres, unsigned int spheres_size, float invWidth, float invHeight, float aspectratio, float angle) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;

    if (y < height && x < width) {
        float xx = (2 * ((x + 0.5) * invWidth) - 1) * angle * aspectratio;
        float yy = (1 - 2 * ((y + 0.5) * invHeight)) * angle;
        Vec3f raydir(xx, yy, -1);
        raydir.normalize();
        image[y*width+x] = trace(Vec3f(0), raydir, spheres, spheres_size, 0);
    }
}

我可以成功编译它:nvcc --ptx kernel.cu -o kernel.ptx ( full PTX here ) 并在带有 cuModuleLoadDataEx 的驱动程序 API 中使用该 PTX使用以下代码段。它按预期工作。

即使我取消注释 #include <math.h> 也能正常工作线(实际上，生成的PTX是完全一样的)。

CudaSafeCall( cuInit(0) );

CUdevice device;
CudaSafeCall( cuDeviceGet(&device, 0) );

CUcontext context;
CudaSafeCall( cuCtxCreate(&context, 0, device) );

unsigned int error_buffer_size = 1024;
std::vector<CUjit_option> options;
std::vector<void*> values;
char* error_log = new char[error_buffer_size];
options.push_back(CU_JIT_ERROR_LOG_BUFFER); //Pointer to a buffer in which to print any log messages that reflect errors
values.push_back(error_log);
options.push_back(CU_JIT_ERROR_LOG_BUFFER_SIZE_BYTES); //Log buffer size in bytes. Log messages will be capped at this size (including null terminator)
values.push_back(&error_buffer_size);
options.push_back(CU_JIT_TARGET_FROM_CUCONTEXT); //Determines the target based on the current attached context (default)
values.push_back(0); //No option value required for CU_JIT_TARGET_FROM_CUCONTEXT

CUmodule module;
CUresult status = cuModuleLoadDataEx(&module, ptxSource, options.size(), options.data(), values.data());
if (error_log && error_log[0]) { //https://stackoverflow.com/a/7970669/3136474
    std::cout << "Compiler error: " << error_log << std::endl;
}
CudaSafeCall( status );

然而，每当我尝试使用 NVRTC ( full PTX here ) 编译这个确切的内核时，它编译成功但给我一个 Segmentation fault (core dumped)正在调用 cuModuleLoadDataEx (当尝试使用生成的 PTX 时)。

如果我取消注释 #include <math.h>行，它在 nvrtcCompileProgram 处失败使用以下输出调用:

nvrtcSafeBuild() failed at cuda_raytracer_nvrtc_api.cpp:221 : NVRTC_ERROR_COMPILATION
Build log:
/usr/include/bits/mathcalls.h(177): error: linkage specification is incompatible with previous "isinf"
__nv_nvrtc_builtin_header.h(126689): here

/usr/include/bits/mathcalls.h(211): error: linkage specification is incompatible with previous "isnan"
__nv_nvrtc_builtin_header.h(126686): here

2 errors detected in the compilation of "kernel.cu".

我用 NVRTC 编译它的代码是:

nvrtcProgram prog;
NvrtcSafeCall( nvrtcCreateProgram(&prog, kernelSource, "kernel.cu", 0, NULL, NULL) );

// https://docs.nvidia.com/cuda/nvrtc/index.html#group__options
std::vector<const char*> compilationOpts;
compilationOpts.push_back("--device-as-default-execution-space");
// NvrtcSafeBuild is a macro which automatically prints nvrtcGetProgramLog if the compilation fails
NvrtcSafeBuild( nvrtcCompileProgram(prog, compilationOpts.size(), compilationOpts.data()), prog );

size_t ptxSize;
NvrtcSafeCall( nvrtcGetPTXSize(prog, &ptxSize) );
char* ptxSource = new char[ptxSize];
NvrtcSafeCall( nvrtcGetPTX(prog, ptxSource) );

NvrtcSafeCall( nvrtcDestroyProgram(&prog) );

然后我只需加载 ptxSource使用前面的代码片段(注意:该代码块与驱动程序 API 版本和 NVRTC 版本使用的相同)。

到目前为止我注意到/尝试过的其他事情

1. PTX generated by the NVCC和 the one generated by NVRTC完全不同，但我无法理解它们以识别可能的问题。
2. 尝试向编译器指定特定的 GPU 架构(在我的例子中是 CC 6.1)，没有区别。
3. 试图禁用任何编译器优化(--ftz=false --prec-sqrt=true --prec-div=true --fmad=false 中的选项 nvrtcCompileProgram)。 PTX 文件变大了，但仍然段错误。
4. 尝试添加 --std=c++11或 --std=c++14到 NVRTC 编译器选项。使用它们中的任何一个，NVRTC 都会生成几乎空的(4 行)PTX，但在我尝试使用它之前不会发出警告或错误。

环境

• SO:Ubuntu 18.04.4 LTS 64 位
• nvcc --version : Cuda 编译工具，版本 10.1，V10.1.168。建立于 Wed_Apr_24_19:10:27_PDT_2019
• gcc --version : gcc (Ubuntu 7.5.0-3ubuntu1~18.04) 7.5.0
• 硬件:英特尔 I7-7700HQ、GeForce GTX 1050 Ti

OP+1 日编辑

我忘了添加我的环境。请参阅上一节。

Also can you compile the nvrtc output with ptxas? – @talonmies' comment

nvcc -生成的 PTX 编译时出现警告:

$ ptxas -o /tmp/temp_ptxas_output.o kernel.ptx
ptxas warning : Stack size for entry function 'raytrace_kernel' cannot be statically determined

这是由于递归内核函数 (more on that)。可以安全地忽略它。

nvrtc -生成的 PTX 不编译并发出错误:

$ ptxas -o /tmp/temp_ptxas_output.o nvrtc_kernel.ptx
ptxas fatal   : Unresolved extern function '_Z5powiffi'

基于 this question我添加了 __device__至 Sphere类构造函数并删除了 --device-as-default-execution-space编译器选项。它现在生成的 PTX 略有不同，但仍会出现相同的错误。

使用 #include <math.h> 编译现在生成很多“没有执行空间注解的函数被认为是宿主函数，宿主函数在JIT模式下是不允许的”。除了先前的错误之外的警告。

如果我尝试使用 accepted solution of the question它抛出一堆语法错误并且无法编译。 NVCC 仍然可以完美运行。

Answer 1

刚刚找到古人的罪魁祸首comment-and-test method : 如果我删除 trace 方法中用于计算菲涅尔效应的 pow 调用，错误就会消失。

现在，我只是将 pow(var, 3) 替换为 var*var*var。

我创建了一个 MVCE并向 NVIDIA 提交错误报告:https://developer.nvidia.com/nvidia_bug/2917596 .

Liam Zhang 回答并指出了我的问题:

The issue in your code is that there is an incorrect option value being passed to cuModuleLoadDataEx. In lines:

options.push_back(CU_JIT_ERROR_LOG_BUFFER_SIZE_BYTES); //Log buffer size in bytes. Log messages will be capped at this size (including null terminator)
values.push_back(&error_buffer_size);

the buffer size option is provided, but instead of passing a value with the size, a pointer to that value is passed. Since this pointer is then read as a number, the driver assumed a much larger buffer size than 1024.

During the NVRTC compilation a "Unresolved extern function" error occurred, because the pow function signature, as you can find in the documentation is:
__device__​ double pow ( double x, double y )
When the driver tried to zero the buffer when putting the error message in it, the segfault happened.
Without the call to pow, there was no compilation error, so the error buffer was not used and there was no segfault.

To ensure the device code is correct, the values used to call pow function as well as the output pointer should be a double number, or a float equivalent function, powf, could be used.

如果我更改对 values.push_back((void*)error_buffer_size); 的调用，它会报告与生成的 PTX 的 ptxas 编译相同的错误:

Compiler error: ptxas fatal   : Unresolved extern function '_Z5powiffi'
cudaSafeCall() failed at file.cpp:74 : CUDA_ERROR_INVALID_PTX - a PTX JIT compilation failed