Callable Programs acting differently from equivalent inlined functions

Gaming and Visualization Technologies Visualization OptiX
1

I’m trying to implement a cosine density PDF, similar to what Peter Shirley does in his “The Rest of Your Life” book. Considering we can’t use virtual functions with OptiX, I tried to use callable programs:

RT_CALLABLE_PROGRAM void cosine_build(onb &uvw, float3 normal){
    uvw.build_from_w(vec3f(normal));
}

RT_CALLABLE_PROGRAM float cosine_value(onb &uvw, float3 direction) {
    float cosine = dot(unit_vector(vec3f(direction)), uvw.w);
    if(cosine > 0.f)
        return cosine / CUDART_PI_F;
    else
        return 0.f;
}

RT_CALLABLE_PROGRAM float3 cosine_generate(onb &uvw, DRand48 &rnd) {
    return uvw.local(random_cosine_direction(rnd)).as_float3();
}

I call these programs from the Raygen program, as follows:

// PDF callable programs
rtDeclareVariable(rtCallableProgramId<void(onb, float3)>, cosine_build, , );
rtDeclareVariable(rtCallableProgramId<float(onb, float3)>, cosine_value, , );
rtDeclareVariable(rtCallableProgramId<float3(onb, DRand48)>, cosine_generate, , );
// ...
onb uvw;
cosine_build(uvw, prd.out.normal.as_float3());
float pdf_val = cosine_value(uvw, prd.out.scattered_direction.as_float3());
float3 pdf_dir = cosine_generate(uvw, rnd);

And assign them in the host side like this:

virtual void assignTo(optix::Context &g_context, optix::Program &program) const override {
        optix::Program build = g_context->createProgramFromPTXString(embedded_cosine_pdf_programs, "cosine_build");
        optix::Program value = g_context->createProgramFromPTXString(embedded_cosine_pdf_programs, "cosine_value");
        optix::Program generate = g_context->createProgramFromPTXString(embedded_cosine_pdf_programs, "cosine_generate");
        
        program["cosine_build"]->setProgramId(build);
        program["cosine_value"]->setProgramId(value);
        program["cosine_generate"]->setProgramId(generate);
    }

where “program” is the Raygen program variable. The issue is, the result is incorrect and completely different than what I expected:

1080×1080 391 Bytes

I tried to copy paste the contents of the callable programs into inlined device functions, like the following, as a test.

inline __device__ void cosine_build(onb &uvw, float3 normal){
  uvw.build_from_w(vec3f(normal));
}

inline __device__ float cosine_value(onb &uvw, float3 direction) {
  float cosine = dot(unit_vector(vec3f(direction)), uvw.w);
  if(cosine > 0.f)
      return cosine / CUDART_PI_F;
  else
      return 0.f;
}

inline __device__ float3 cosine_generate(onb &uvw, DRand48 &rnd) {
  return uvw.local(random_cosine_direction(rnd)).as_float3();
}

I didn’t change how these functions are called from the raygen program, or anything else in the project. To my surprise, it worked as initially expected:

1080×1080 804 KB

My question is, why exactly am I seeing this behavior? The functions’ content is fundamentally the same as the callable programs, so why exactly is the result completely different?

I’m currently using CUDA 10 and OptiX 5.1 on a Windows 10 machine.

2

It seems, your programs are not properly assigned.

program["cosine_build"]->setProgramId(build);
program["cosine_value"]->setProgramId(value);
program["cosine_generate"]->setProgramId(generate);

callable programs are also variables, they need to be assigned to a proper scope. For example, you can assign them to context:

g_context["cosine_build"]->setProgramId(build);
g_context["cosine_value"]->setProgramId(value);
g_context["cosine_generate"]->setProgramId(generate);
3

I’m not sure if that is the issue, it looks like they were already being bound to some program in their example.

Could it be with the callable program variable defenition? I don’t think your passing onb as a reference in both the callable program function AND the OptiX variable definition.

The perRayData/State/Paramaters in the advanced samples is passed in as such.

rtBuffer< rtCallableProgramId<void(MaterialParameter const& parameters, State const& state, PerRayData& prd)> > sysSampleBSDF;
4

I just checked optixCallablePrograms.cpp. Callable programs can also be set to the scope of a program:

Program exception_program = context->createProgramFromPTXString( ptx, "exception" );
  context->setExceptionProgram( 0, exception_program );
  context["bad_color"]->setFloat( 1.0f, 1.0f, 0.0f );

  // Miss program
  ptx = sutil::getPtxString( SAMPLE_NAME, "optixCallablePrograms.cu" );
  Program miss_program  = context->createProgramFromPTXString( ptx, "miss" );
  Program shade_program = context->createProgramFromPTXString( ptx, "shade_from_ray" );

  Buffer func_buffer = context->createBuffer( RT_BUFFER_INPUT, RT_FORMAT_PROGRAM_ID, 1 );
  int* func_data = static_cast<int*>( func_buffer->map() );
  func_data[0] = shade_program->getId();
  func_buffer->unmap();

  miss_program[ "shade_ray" ]->set( func_buffer );
  context->setMissProgram( 0, miss_program );
5

Radon nailed it. Despite the fact I was using references in the program definitions, that wasn’t the case in the rtDeclareVariables, which means I was just passing by value again and again.

This fixed it:

rtDeclareVariable(rtCallableProgramId<void(onb&, float3)>, cosine_build, , );
rtDeclareVariable(rtCallableProgramId<float(onb&, float3)>, cosine_value, , );
rtDeclareVariable(rtCallableProgramId<float3(onb&, DRand48&)>, cosine_generate, , );

Thanks a lot for the help, Radon and yashiz!