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

Question

总结

我正在将一个基于 Scratchapixel version 的简单光线追踪应用程序移植到一堆 GPU 库中。我使用运行时 API 和驱动程序 API 成功地将它移植到 CUDA，但是当我尝试使用运行时编译的 PTX 和 NVRTC 时它抛出 Segmentation fault (core dumped)。如果我取消注释内核文件开头的 #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) 并在驱动程序 API 和 cuModuleLoadDataEx 中使用该 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]) { //
    std::cout << "Compiler error: " << error_log << std::endl;
}
CudaSafeCall( status );

然而，每当我尝试使用 NVRTC (full PTX here) 编译这个确切的内核时，它编译成功但在调用 cuModuleLoadDataEx 时给了我一个 Segmentation fault (core dumped) （当试图使用生成的 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 版本使用的相同）。

到目前为止我 noticed/tried 的其他内容

PTX generated by the NVCC and the one generated by NVRTC 完全不同，但我无法理解它们以找出可能的问题。
尝试向编译器指定特定的 GPU 架构（在我的例子中是 CC 6.1），没有区别。
试图禁用任何编译器优化（nvrtcCompileProgram 中的选项 --ftz=false --prec-sqrt=true --prec-div=true --fmad=false）。 PTX 文件变大了，但仍然 Segfaulting.
尝试将 --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
硬件：Intel 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'

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

使用 #include <math.h> 编译现在除了以前的错误之外还会生成很多 "A function without execution space annotations is considered a host function, and host functions are not allowed in JIT mode." 警告。

如果我尝试使用它会抛出一堆语法错误并且无法编译。 NVCC 仍然完美运行。

Answer 1

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

目前，我刚刚将 pow(var, 3) 替换为 var*var*var。

我创建了一个MVCE and filled a bug report to 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

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

Differences between NVCC and NVRTC on compilation to PTX

c++

cuda

ptx

nvrtc

总结

详细说明

到目前为止我 noticed/tried 的其他内容

环境

OP+1 日编辑