Failed to translate PTX input to LLVM

Hi all,
I am receiving this error:

error occurred with error code 1287 and message Parse error (Details: Function "RTresult _rtProgramCreateFromPTXString(RTcontext, const char*, const char*, RTprogram_api**)" caught exception: (api input string): error: Failed to translate PTX input to LLVM
Unsupported array initializer elem type4
PTX Symbols (e.g. pointers to functions) cannot be used to initialize arrays.

when I try to create a closest hit program.
Is there any way to know which symbol/array or line code it is referring to, in order to debug the program?

I can provide the PTX string if necessary or any other information.

I am using Optix 6.0 with CUDA 10.1 on Ubuntu 18.04 with gcc 7.3.0.

Thanks a lot

Please always read the OptiX Release Notes before setting up a development environment for a new OptiX version.

You did it again. CUDA 10.1 is not on the list of supported CUDA toolkits inside the OptiX 6.0.0 Release Notes. CUDA 10.0 is the maximum tested version.

If it fails with a supported CUDA version as well, what have you done to isolate the error?
What are the nvcc compiler options? Have you tried the SM 3.0 architecture as target?
Is there any array initialization inside your CUDA source code?

Hi again,
I am sorry about my lack of care. I have recompiled and tried all the code with the supported version 10.0.130. Here is my info:

[1][SYS INFO    ]
OptiX Version:[6.0.0] Branch:[r418_00] Build Number:[25823979] CUDA Version:[cuda100] 64-bit
Display driver: 418.43
Devices available:
CUDA device: 0
    0000:01:00.0
    GeForce GTX 970
    SM count: 13
    SM arch: 52
    SM clock: 1240 KHz
    GPU memory: 4040 MB
    TCC driver: 0
    Compatible devices: 0

I use the sutil library provided by the SDK. The NVRTC options it uses (printed from sutil.cpp) are:

-I/home/eegea/optix6/SDK/multitransmitter
-I/home/eegea/optix6/include
-I/home/eegea/optix6/include/optixu
-I/home/eegea/optix6/SDK/support/mdl-sdk/include
-I/usr/local/cuda-10.0/include
-I/home/eegea/optix6/SDK/sutil
-I/home/eegea/optix6/SDK/cuda
-arch
compute_30
-use_fast_math
-lineinfo
-default-device
-rdc
true
-D__x86_64

Apparently the problem is in this line in the CU file I use for the program:

HitInfoMultiTransmitter internalHit;

where the HitInfoMultiTransmitter is a struct defined as:

struct HitInfoMultiTransmitter  {
        optix::uint4 whrd; // [written,refhash,rxBufferIndex,dist] 
        optix::float2 E;   // Complex
        optix::uint transmitterIndex;

         __device__ HitInfoMultiTransmitter() : whrd(make_uint4(0u,0u,0u,0u)),E(make_float2(0.0f,0.0f)),transmitterIndex(0)
        {
        }
        //Equality operator: hits are equal if the combined has is equal, that is, the have hit the same sequence of faces
         __device__ bool operator==(const HitInfoMultiTransmitter &h) const {
                return (whrd.y == h.whrd.y);
        };
        //Sorting. Here we first order by txId (tx buffer index), then check receiver id (receiver buffer index), then hash and finally distance. Hits are ordered by txId, rxid, combined_hash and distance to receiver
         __device__ virtual bool operator<(const HitInfoMultiTransmitter &h) const {
                if (transmitterIndex==h.transmitterIndex){
                        if (whrd.z == h.whrd.z) {
                                if (whrd.y == h.whrd.y) {
                                        return (whrd.w < h.whrd.w);
                                } else {
                                        return (whrd.y < h.whrd.y);
                                }
                        } else {
                                return (whrd.z < h.whrd.z);
                        }
                } else {
                        return (transmitterIndex<h.transmitterIndex);
                }
        };
};

If I comment out the line where this struct is used the program is compiled and the launch proceeds, otherwise, it does not allow me to generate it.
I have tried different ways to initialize the struct but none seems to work.

I need that struct because I use thrust to process on device some results.

A previous program very similar but that uses this slightly different struct works correctly without any problem:

struct HitInfo {
        //Packed to fullfill alignment rules, see for instance https://devtalk.nvidia.com/default/topic/1037798/optix/optix-time-for-launch/
        optix::uint4 whrd; // [written,refhash,rxBufferIndex,dist] 
        optix::float2 E;   // Complex

        //Equality operator: hits are equal if the combined has is equal, that is, the have hit the same sequence of faces
        __host__ __device__ bool operator==(const HitInfo &h) const {
                return (whrd.y == h.whrd.y);
        };
        //Sorting. First check id (buffer index), then hash and finally distance. Hits are ordered by id, combined_hash and distance to receiver
        __host__ __device__ bool operator<(const HitInfo &h) const {
                if (whrd.z == h.whrd.z) {
                        if (whrd.y == h.whrd.y) {
                                return (whrd.w < h.whrd.w);
                        } else {
                                return (whrd.y < h.whrd.y);
                        }
                } else {
                        return (whrd.z < h.whrd.z);
                }
        };
}

Thank you very much for your help

Try removing the “virtual” from the non-working version.
device virtual bool operator<(const HitInfoMultiTransmitter &h) const

Dynamic function (tables) in OptiX need to be implemented with (buffers of) bindless callable program IDs.

I would also recommend to use forceinline on all device functions.
Every device function which is not RT_PROGRAM or RT_CALLABLE_PROGRAM in OptiX device code should be prefixed with forceinline.
I’m using this RT_FUNCTION define to make code look more consistent:
https://github.com/nvpro-samples/optix_advanced_samples/blob/master/src/optixIntroduction/optixIntro_07/shaders/rt_function.h

Yes, that solved the problem.
Thanks a lot for your help.
Kind regards