Undefined behavior in device code

Gaming and Visualization Technologies Visualization OptiX
6

Hi,

I have been able to narrow down the problem to an as minimal as possible reproducer. To reproduce this issue, create a file called tmp.h in the optixPathTracer directory with the contents: pastebin. Prepend the optixPathTracer.h file with pastebin and the following line to the end of the Params struct in optixPathTracer.h:

s3 *d_s3is;

Now, add the following lines to the top of the launchSubframe method of optixPathTracer.cpp:

CUDA_CHECK(cudaMalloc(
        reinterpret_cast<void **>(&state.params.d_s3is),
        sizeof(s3)));
s3 s3i0;
s3i0.v6.v2 = false;
s3i0.v6.v3 = true;
CUDA_CHECK(cudaMemcpy(
        reinterpret_cast<void *>(state.params.d_s3is),
        &s3i0,
        sizeof(s3), cudaMemcpyHostToDevice));

and the following line to the end of the same method:

CUDA_CHECK(cudaFree(reinterpret_cast<void *>(state.params.d_s3is)));

Finally, add the line:

eval();

to the top of the ray gen method (__raygen__rg) in optixPathTracer.cu and replace in the same file

#include <optix.h>

#include "optixPathTracer.h"
#include "random.h"

#include <sutil/vec_math.h>
#include <cuda/helpers.h>

extern "C" {
__constant__ Params params;
}

by:

#include <optix.h>

#include "tmp.h"
#include "random.h"

#include <sutil/vec_math.h>
#include <cuda/helpers.h>

(note that the params are removed, I define these in tmp.h) If you now run optixPathTracer (I have done so with a CMake “Release” build), the result should be a failed assertion and a crash as a result. If you un-comment the assertion on line 56 of tmp.h, the program should run fine. I have annotated the problem with comments the best I could. I these instructions for reproducing are unclear (or anything else about the problem), please let me know. I would have uploaded the full sample on pastebin but I was not sure if distributing the sample code is allowed.

To also respond to your remarks:

Yes I am. I thought that this may be a CUDA-related issue, but from what I understand, OptiX 7.3 only works with CUDA 11.1 so I cannot upgrade the CUDA version.

There is no optixReportIntersection in the original intersection program. I just directly perform calculations in the intersection program. That pipeline has no any- or closest-hit programs.