文章/答案/技术大牛

发布

社区首页 >问答首页 >NVCC与NVRTC在PTX编译上的差异

问NVCC与NVRTC在PTX编译上的差异
EN

Stack Overflow用户

提问于 2020-04-01 03:37:29

回答 1查看 2.2K关注 0票数 2

摘要

我正在将一个基于Scratcha像素版本的简单光线跟踪应用程序移植到一堆GPU库中。我成功地使用运行时API和驱动程序API将其移植到CUDA，但是当我尝试使用在运行时使用NVRTC编译的PTX时，它会抛出一个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 (全PTX在这里)成功地编译它，并使用下面的代码片段在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 (全PTX在这里)编译这个精确的内核时，它都会成功编译，但是在调用cuModuleLoadDataEx时(当尝试使用生成的PTX时)，它会给我一个Segmentation fault (core dumped)。

如果取消对#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版本使用的代码块相同)。

我注意到/尝试过的其他事情

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

环境

因此:Ubuntu18.04.4 LTS 64位
nvcc --version：Cuda编译工具，版本10.1，V10.1.168。建立在Wed_Apr_24_19:10:27_PDT_2019上
gcc --version：gcc (Ubuntu7.5.0-3 ubuntu1~18.04) 7.5.0
硬件: Intel I7-7700HQ，GeForce GTX 1050 Ti

在OP+1日编辑

我忘了增加我的环境。见前一节。

还能用ptxas编译nvrtc输出吗？-@talonmies的注释

nvcc-generated PTX编译时有一个警告：

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

这是由于递归内核函数(更多关于这一点)。它可以安全地被忽略。

nvrtc-generated PTX执行而不是编译，并发出错误：

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

基于这个问题，我将__device__添加到Sphere类构造函数中，并删除了--device-as-default-execution-space编译器选项。它现在生成一个稍微不同的PTX，但仍然显示相同的错误。

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

如果我尝试使用已接受的问题解决办法，它会抛出大量语法错误而不编译。NVCC的工作仍然完美无缺。

nvrtc

c++

cuda

ptx

回答 1

Stack Overflow用户

回答已采纳

发布于 2020-04-02 01:37:00

刚刚发现了古代评测法的罪魁祸首:如果删除用于计算trace方法中的菲涅尔效应的pow调用，错误就会消失。

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

我创建了一个MVCE，并向NVIDIA：bug/2917596填写了一个bug报告。

利亚姆·张回答并指出了这个问题：

代码中的问题是，向cuModuleLoadDataEx传递的选项值不正确。行： options.push_back(CU_JIT_ERROR_LOG_BUFFER_SIZE_BYTES)；//日志缓冲区大小(以字节为单位)。日志消息将被限制在此大小(包括空终止符) values.push_back(&error_buffer_size)；提供了缓冲区大小选项，但不是传递带有大小的值，而是传递指向该值的指针。由于该指针随后被读取为一个数字，驱动程序假定的缓冲区大小比1024大得多。在NVRTC编译过程中，出现了一个“未解决的extern函数”错误，因为pow函数签名，如您在文档中所发现的： __device__ double pow ( double x, double y ) 当驱动程序在输入错误消息时试图使缓冲区为零时，就会发生分段错误。如果没有对pow的调用，就没有编译错误，因此没有使用错误缓冲区，也没有分段错误。为了确保设备代码是正确的，可以使用用于调用pow函数以及输出指针的值为双数字，或者使用浮点等效函数powf。

如果我将调用更改为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

票数 1

页面原文内容由Stack Overflow提供。腾讯云小微IT领域专用引擎提供翻译支持

原文链接：

https://stackoverflow.com/questions/60963315

复制

相似问题

问NVCC与NVRTC在PTX编译上的差异
EN

回答 1

Stack Overflow用户

社区

活动

圈层

关于

腾讯云开发者

热门产品

热门推荐

更多推荐

问NVCC与NVRTC在PTX编译上的差异EN

回答 1

Stack Overflow用户

社区

活动

圈层

关于

腾讯云开发者

热门产品

热门推荐

更多推荐

问NVCC与NVRTC在PTX编译上的差异
EN