OptiX 6.5 CUDA 11.6 Unknown usage of symbol: __cudart_i2opi_d

I believe that I have tracked this down to the use of double precision trig functions, cos/sin/etc because the error goes away if I remove the use of sin(double) and cos(double) although, I do not know or understand how to resolve it without removing the double precision functions.

This same code is currently working without issue on CUDA 11.1

The full error when attempting to load the PTX file is

Unknown error (Details: Function "RTresult _rtProgramCreateFromPTXString(RTcontext, const char*, const char*, RTprogram_api**)" caught exception: Compile Error: Unknown usage of symbol: __cudart_i2opi_d at: [ Instruction:   %11 = extractelement <2 x i32> %10, i32 0, contained in basic block: $L__BB1_3, in function: __internal_trig_reduction_slowpathd, in module: Canonical__Z14ray_generationv from (api input string) ])

The working environment is; Windows 10, OptiX 6.5, Driver Version: 456.71, CUDA Version: 11.1 with a Quadro P3200
A secondary working environment is; Redhat8 with a V100 (I don’t have access to the other details, but it is not CUDA 11.6)

The non-working environment is; Redhat8, OptiX 6.5, Driver Version: 510.39.01, CUDA Version: 11.6 with a Quadro RTX 5000

Any assistance with this would be greatly appreciated.

-Scott

Thanks for the report. Do you have a reproducer for this you can share? Or, if not, would it be possible to send the input PTX? (You can DM me privately if you prefer.)

Are you able to work around this on your non-working system by using either an older driver or an older CUDA toolkit?


David.

This is part of a much larger code base, but today I added a #define USE_DOUBLE_TRIG 0 and just swapped any calls I made to cos(double) with cosf(float) and everything works as expected. I haven’t tried, but highly suspect that if you literally had a ray_generation program with a sin or cos that you can repro.

The systems are administered by a larger group so I do not have the ability to downgrade (very easily) but, as I indicated above, I am just disabling double precision trig for now. I really want to understand how what is happening because I really need to leverage doubles.

// commonly used trig function that causes the reported error on CUDA 11.6
static __device__ __inline__ double sind(const double angle,
                                         const double epsilion = 1e-14)
{
#if USE_DOUBLE_TRIG
    const double result = sin(angle);
    return fabs(result) < epsilion ? 0.0 : result;
#else
    const float result = sinf(static_cast<float>(angle));
    return fabsf(result) < epsilion ? 0.0 : result;
#endif
}

Okay, I’m feeling pretty dumb here. I do not see a way to initiate a PM.

Message them by selecting their avatar and using the message button

Instructions unclear; no message button found!

If you’re logged in and click on the round green NVIDIA logo next to David’s post, there should pop up a window with a big green Message button on the right.
Let us know if not, and we’ll inform the forum team.
(That Message button does not appear if you click on your own avatar image.)

I do not see it on Chrome or FireFox

If you click on the “+7 More” button, it’ll take you to the user page, and there should be a “Message” button at the top.


David.

Still not seeing it. Must be that double precision trig.

This thread is :fire:! Love it.

:joy: That’s fun, not sure why the message buttons are missing. I PM’d you, and saw the acknowledgement. So for anyone else reading, if you can’t find the message button, feel free to request an initial message from one of us.


David.

I did modify the optixTutorial sample to use double precision sin(), and on Windows with CUDA 11.6 it didn’t cause me any trouble yet. See if you can see the same thing. There’s a sin() function in tutorial5.cu, which you can invoke using “-T 5”. I also put a double precision sin() call in the raygen program (called “pinhole_camera”). I dropped a printf in there in both cases to verify things we recompiled and called.

For posterity, the missing message buttons are an intentional feature of this forum software for brand new user accounts to prevent spam.


David.

1 Like

I do not have CUDA 11.6 on a Windows machine, but I certainly trust that you were successful with that. My problem platform is RHEL8 with gcc10 and a Quadro RTX 5000 Driver Version: 510.39.01. Another difficult aspect is that I am running remotely with a terminal only. Fear Not!

I have been able to reproduce this issue by modifying the optixConsole sample provided by OptiX 6.5. I made the following minimum changes and generated the issue.

  1. Modified the main CMakeLists.txt file line 154 changing the shader model from sm_30 to sm_70 because I got a build error about no support for sm_30.
  2. Modified optixContext/constantbg.cu by grabbing some trig code from optixTutorial/tutorial5.cu envmap_miss() see below: with the #if 1 it works and gives my a bit of a gradient background, setting #if 0 causes the issue.
rtDeclareVariable(optix::Ray, ray,          rtCurrentRay, );
rtDeclareVariable(PerRayData_radiance, prd_radiance, rtPayload, );

RT_PROGRAM void miss()
{
#if 0
  float theta = atan2f( ray.direction.x, ray.direction.z );
  float phi   = M_PIf * 0.5f -  acosf( ray.direction.y );
  float u     = (theta + M_PIf) * (0.5f * M_1_PIf);
  float v     = 0.5f * ( 1.0f + sin(phi) );
#else
  double theta = atan2( (double)ray.direction.x, (double)ray.direction.z );
  double phi   = M_PIf * 0.5 - acos( (double)ray.direction.y );
  double u     = (theta + M_PIf) * (0.5 * M_1_PIf);
  double v     = 0.5 * ( 1.0 + sin(phi) );
#endif

  //prd_radiance.result = bg_color;
  prd_radiance.result = make_float3(0.0f, u, v);
}

The above should be enough to confirm on your end. I would also suspect that your tutorial5.cu change would also fail in the same way on a RHEL8 machine. It looks like the samples are using the Runtime Compiler (no ptx) so I can just change that #if and re-run the sample.

Thanks

Thank you for the reproducer! This does reproduce on my Ubuntu 20 machine, and your reproducer has been filed with the bug report. I will update when the fix is scheduled for release. From my testing, this does indeed appear to be introduced by a CUDA toolkit, somewhere between 11.1 (no repro) and 11.4 (repro). Behavior does not appear to change based on driver (even though I used the OptiX default NVRTC run-time compilation setup).

I believe what this means is that the fix will need to come in the form of an updated CUDA toolkit. You can see the pace of the recent minor version releases: https://developer.nvidia.com/cuda-toolkit-archive Unless we can find a workaround, I will recommend sticking on CUDA 11.1 in order to use double precision until a toolkit with the fix is released.


David.

1 Like