Hello, every one!
I am trying to build multiple modules with multiple .ptx files with Optix 7.1, because I want to split my code into several .cu files for better code management.
And I noticed that if there is no Preformatted text __raygen__ function in the .cu file (which is compiled into a .ptx file later), optixModuleCreateFromPTX() will return an 7200 error.
I tried this because I want to move my __raygen__ function into a separate file. Currently I solve this problem by adding a dummy __raygen__ function into .cu files which don’t have one.
So my question is:
Is there any better way to avoid the 7200 error rather than create dummy functions?
What is the purpose to test if there is a __raygen__ function in .ptx file, while you are able to create multiple modules and choose functions from any of them?
Thank you.
It’s not exactly a __raygen__ which is required, but yes, OptiX does not handle modules without at least one of the program domains defined inside each module.
Please see this thread for a similar workaround using a unique dummy __direct_callable__
https://forums.developer.nvidia.com/t/optix-7-equivalent-of-rtprogramcreatefromptxfiles/129321/8
I’m also using multiple modules in my examples, but they are grouped into the main program domains, so that there is alway at least one of the semantic programs inside each module:
https://github.com/NVIDIA/OptiX_Apps/blob/master/apps/rtigo3/src/Device.cpp#L528
For performance reasons I do not split regular functions into separate modules which are then linked.
OptiX wants code to be inlined as much as possible. Because of that, I define functions used in multiple modules as __forceinline__ __host__ __device__ inside headers.
Examples here:
https://github.com/NVIDIA/OptiX_Apps/blob/master/apps/rtigo3/shaders/vector_math.h
https://github.com/NVIDIA/OptiX_Apps/blob/master/apps/rtigo3/shaders/shader_common.h
Sorry but can you explain what is a program domain?
And how to group modules into one domain?
With program domains I just mean the OptiX program types in the table of this chapter:
https://raytracing-docs.nvidia.com/optix7/guide/index.html#program_pipeline_creation#program-input
In the Device.cpp code link above you can see that I have raygen, miss, exception, closest hit and any hit programs as well as different direct callable programs in different PTX files. That’s all.
None of these hit your initial issue because there is always at last one of the program types inside each PTX with this code structure.
That partitioning into separate files isn’t necessary and you could also put everything into one source file.
It’s simply a matter of balancing compile times during development and initial startup times during runtime.
Since OptiX caches the internally assembled code and CUDA does the same with the final microcode, it’s only the initial startup of the application where the pipeline creation path is slow.
You can see that when enabling the OptixDeviceContextOptions callback at level 4, which will print out the program cache misses and hits.
Find more information about both of these topics in this chapter of the programming guide:
https://raytracing-docs.nvidia.com/optix7/guide/index.html#context#context
Thank you so much for your replies.
In fact I find my case is kind of complex after I did some experiments .
Before the separation, my only .cu file contained every Optix programs like this:
https://github.com/gzfrozen/EM_tracing/blob/master/src/cuda/devicePrograms.cu#L73
Then I tried to separate the .cu file, first step by only pick out the __raygen__ program, separate it into environmentRender.cu and phaseDetection.cu ( rayLaunch.cu is not been used by now) as following github branch:
However in this case, despite the original big .cu file (environmentRender.cu) still contains ton’s of Optix programs, an extra dummy __raygen__ program is needed or the module can not be created and return this:
[ 2][COMPILE FEEDBACK]: COMPILE ERROR: Invalid PTX input: ptx2llvm-module-001: error: Failed to parse input PTX string
ptx2llvm-module-001, line 9; warning : Unsupported .version 6.5; current version is ‘6.4’
ptx2llvm-module-001, line 460; fatal : Parsing error near ‘.version 6.5’: syntax error
Cannot parse input PTX string
Optix call (optixModuleCreateFromPTX(optixContext, &moduleCompileOptions, &pipelineCompileOptions, ptx.c_str(), ptx.size(), log, &sizeof_log, &m)) failed with code 7200 (line 377)
Magical things happened after I tried further separation into 3 .cu files environmentRender.cu closestHit.cu rayLaunch.cu in this branch:
Suddenly the dummy __raygen__ is no longer needed and can be safely commented out.
So I guess there might be a bug in module compiling?
Anyway, thank you again for giving me the advice to separate .cu files. I think I will try to separate the .cu files by program types like this example you have shown me:
https://github.com/NVIDIA/OptiX_Apps/blob/master/apps/rtigo3/src/Device.cpp#L528