Custom Intersection Program - OptiX 7

Professional Graphics and Rendering Advanced Graphics OptiX
#1

Hello all,

I apologize if this is too simple a question, but here goes. Is it possible to create a custom intersection program with OptiX 7 such that an integer parameter is passed to it? for example, I have a .cu file with the following CUDA code snippet:
extern "C" __global__ void __intersection__quad(int qIdx) { ... }
which is successfully compiled to PTX code, which for purposes of this question I will call is_ptx (a string). Later, in my main.cpp file I go to build an OptixModule from the PTX code as follows:
OptixPipelineCompileOptions pipeline_compile_options = {};

  // Default options MUST be consistent for all modules used in a SINGLE pipeline
  OptixModuleCompileOptions module_compile_options = {};
  module_compile_options.maxRegisterCount     = OPTIX_COMPILE_DEFAULT_MAX_REGISTER_COUNT;
  module_compile_options.optLevel             = OPTIX_COMPILE_OPTIMIZATION_DEFAULT;
  module_compile_options.debugLevel           = OPTIX_COMPILE_DEBUG_LEVEL_LINEINFO;
  pipeline_compile_options.usesMotionBlur     = false;

  // This option is important to ensure we compile code which is optimal for 
  // out scene hierarchy. We use a SINGLE GAS - no instancing or multi-level
  // hierarchies
  pipeline_compile_options.traversableGraphFlags = OPTIX_TRAVERSABLE_GRAPH_FLAG_ALLOW_SINGLE_GAS;

  // Our DEVICE code uses 2 payload registers
  pipeline_compile_options.numPayloadValues      = 2;
  pipeline_compile_options.numAttributeValues    = 2;
  pipeline_compile_options.exceptionFlags = OPTIX_EXCEPTION_FLAG_STACK_OVERFLOW |
                                            OPTIX_EXCEPTION_FLAG_TRACE_DEPTH |
                                            OPTIX_EXCEPTION_FLAG_USER |
                                            OPTIX_EXCEPTION_FLAG_DEBUG;

  // This is the name of the launch parameter struct in our DEVICE code
  // e.g. LaunchParameter params
  pipeline_compile_options.pipelineLaunchParamsVariableName = "params";

OptixModule moduleIS = nullptr;
  {
    // Get PTX shader IntersectionReflectQuad - i.e. Intersection program
    std::cout << "Reading PTX shader (Intersection) that was previously compiled ...\n";
    std::string is_ptx;
    getShaderString(optix7_ptx::IntersectReflectQuad, is_ptx);
  
    std::cout << "Creating Intersection module ...\n";
  
    // Create actual module from PTX
    OPTIX_CHECK_LOG(optixModuleCreateFromPTX(context, &module_compile_options,
                                             &pipeline_compile_options, is_ptx.c_str(),
                                             is_ptx.size(), log, &sizeof_log, &moduleIS));
  }

Where ‘context’ above is GPU DEVICE context, successfully set previously in file. However, when I try to run this I get a segmentation fault.

I am guessing that maybe there needs to be a special setting in compile/pipeline options - or maybe I am completely off the mark. Hopefully something very simple :)

Anyway, any hints/help would be greatly appreciated.

Thanks

#2

Hey there,

So the normal mechanisms to pass values to your intersection program are via either your payload, or your shader binding table entry. You would normally use the payload to communicate between raygen or a hit-shader and the intersection shader. Since the payload uses registers, it will be faster than storing parameters in memory.

The payload values are 32 bit integers, you actually have to do more work to use any other type, like float. There are no compile options or pipeline options to enable different types, it’s just up to your code to interpret the payload bits consistently on both ends. For example, to use floats typically you call float_as_uint() and match that with a call to uint_as_float() on the other side. This is true for both payload values and attribute values.

https://raytracing-docs.nvidia.com/optix7/guide/index.html#device_side_functions#payload-access

While not recommended unless you have no other choice, if you pass your parameters via an SBT entry pointer to memory, then you will have defined the types yourself, so there is no type-casting to worry about.

If you compiled an intersection program that takes a parameter, that may be the cause of your segmentation fault, since OptiX does not pass a parameter to this function directly using typical function calling mechanisms. If you do not have a parameter in your compiled intersection program, then the segmentation fault means you’ve tried to access memory that is out of bounds.

–
David.

#3

Thanks @dhart for the reply. A lot of useful information. I thought there might be an issue with a compiled intersection program that takes a parameter, but didn’t find it until I tried to build the module. Learned something new today.

Do you have a link to any code that shows an actual __intersection__ program in OptiX 7?

Thank you again for your help.

#4

Sure, in the 7.2 SDK you can find intersection programs in optixWhitted/geometry.cu, optixDemandTexture/optixDemandTexture.cu, optixDynamicMaterials/optixDynamicMaterials.cu, optixSimpleMotionBlur/optixSimpleMotionBlur.cu.

It appears that all of them use the SBT to pass some parameters and none of them have an example of using the payload. This is because they need per-primitive data that is set by the host-side program, so in that case there’s no avoiding a read from memory.

If you want to use the payload to pass values from raygen or a hit program to your intersection program, all you need to do is read the payload values using optixGetPayload_n(), where n is the numeric payload register to read, and it will return whatever value you passed into optixTrace() for that register.

–
David.

#5

Cool - Thank you.

My original idea of using the intersection program was to pass an index as an argument to the __intersection__ program which is obviously not possible in OptiX 7. This index was to an index into quads.

So I was thinking of maybe passing that index value via an associated SBT, however I am not 100% on how the SBT relates to a given shader/program (e.g. __intersection__ ). Can you tell me how a given SBT relates to a specific shader/program? Is there a single SBT for all Rays in a given hit group (like the __intersection__ program)? Can multiple SBT(s) be assigned to the same program?

Thank you again for the assist and I apologize for the rambling.

#6

Tell me more about how you want to use this index, I’m not sure I understand the goal. Do you just need to get the vertices of this quad so you can compute the intersection? Or are you trying to do something fancier like controlling layers, or have all rays intersect one primitive at a time, or something else?

If you have a set of quads in your scene, and you just need to identify which quad your ray is intersecting so that you can load its vertices and compute the intersection, then you could use optixGetPrimitiveIndex() rather than passing the index to your intersector. The value from optixGetPrimitiveIndex() will match whatever order you used for your bounding boxes.

The SBT, or Shader Binding Table, is the array that OptiX uses to figure out which shader to call for each interaction between a ray and your geometry. It’s nothing more than a lookup table that lets you have separate shaders for every combination of geometry type and material type and ray type, and the result of the lookup is to call the shader at the entry in your table.

There is one SBT per launch, and the SBT contains entries for 1 raygen program, one optional miss program per ray type, and then optional intersection, any-hit and closest-hit programs for the different geometry / material / ray types.

You have a lot of control over how you arrange your SBT, you can share the same programs across multiple material or geometry or ray types. In OptiX, you have the option to attach your own arbitrary data to each entry in the table. A lot of the apparent complexity of the SBT just comes from it’s flexibility, but the basics are pretty straightforward.

https://raytracing-docs.nvidia.com/optix7/guide/index.html#shader_binding_table#shader-binding-table

https://www.willusher.io/graphics/2019/11/20/the-sbt-three-ways

–
David.

1 Like
#7

Thank you @dhart for the detailed and professional response - much appreciated.

I will be using the index in the latter form (nothing fancy, just vertices of the quad), therefore it looks like optixGetPrimitiveIndex( ) will be what I will use.

The SBT do seem complex but when you break it down as a lookup table for program(s) it makes more sense. Thank you again for the information, most helpful.

1 Like