Module creation not identifying my __direct_callable__ functions

Hi all,

I’m in the process of moving my lens shaders into direct callable functions as in the advanced examples Optix_Apps Lens Shaders. But when I try to create my module I get an error that it is invalid. If I add for example a raygen program to the same ptx the raygen program is identified but not the direct callables.

My project is originally based on Ingo Walds Siggraph course so I’m using that compile and embed macro to create the ptx. Not sure but maybe the compiler can’t identify the programs?

macro(cuda_compile_and_embed output_var cuda_file)
  set(c_var_name ${output_var})
  cuda_compile_ptx(ptx_files ${cuda_file})
  list(GET ptx_files 0 ptx_file)
  set(embedded_file ${ptx_file}_embedded.c)
#  message("adding rule to compile and embed ${cuda_file} to \"const char ${var_name}[];\"")
  add_custom_command(
    OUTPUT ${embedded_file}
    COMMAND ${BIN2C} -c --padd 0 --type char --name ${c_var_name} ${ptx_file} > ${embedded_file}
    DEPENDS ${ptx_file}
    COMMENT "compiling (and embedding ptx from) ${cuda_file}"
    )

  set(${output_var} ${embedded_file})
endmacro()

include_directories(${OptiX_INCLUDE})
add_definitions(-D__CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__=1)

Also here’s an example of my direct callable:

  extern "C" __device__ void __direct_callable__rectilinear(const uint3 launchIndex, const uint3 launchDim, vec3f & position, vec3f & rayDir)
  {
	  const auto& viewpoint = params.viewpoint;
	  const auto& camera = params.camera;

	  const vec2f screen(vec2f(launchIndex.x + .5f, launchIndex.y + .5f) / vec2f(launchDim.x, launchDim.y));

	  // generate ray direction
	  rayDir = normalize(viewpoint.direction * camera.FoVdepth
		  + (screen.x - 0.5f) * viewpoint.horizontal
		  + (0.5f - screen.y) * viewpoint.vertical * camera.aspect);

	  position = viewpoint.position;
  }

and my module creation:

OptixModuleCompileOptions   moduleCompileOptions = {};

moduleCompileOptions.maxRegisterCount  = 0;
moduleCompileOptions.optLevel          = OPTIX_COMPILE_OPTIMIZATION_DEFAULT;
moduleCompileOptions.debugLevel        = OPTIX_COMPILE_DEBUG_LEVEL_NONE;

pipelineCompileOptions.traversableGraphFlags			= OPTIX_TRAVERSABLE_GRAPH_FLAG_ALLOW_SINGLE_LEVEL_INSTANCING;
pipelineCompileOptions.usesMotionBlur					= false;
pipelineCompileOptions.numPayloadValues				= 5;
pipelineCompileOptions.numAttributeValues				= 2;
pipelineCompileOptions.exceptionFlags					= OPTIX_EXCEPTION_FLAG_NONE; //OPTIX_EXCEPTION_FLAG_DEBUG;
pipelineCompileOptions.pipelineLaunchParamsVariableName = "params";
pipelineCompileOptions.usesPrimitiveTypeFlags			= OPTIX_PRIMITIVE_TYPE_FLAGS_TRIANGLE;

const std::string ptxCodeLensShaders = embedded_ptx_lensShaders;

char log[2048];
size_t sizeof_log = sizeof(log);
OPTIX_CHECK(optixModuleCreateFromPTX(optixContext,
	&moduleCompileOptions,
	&pipelineCompileOptions,
	ptxCodeLensShaders.c_str(),
	ptxCodeLensShaders.size(),
	log, &sizeof_log,
	&moduleDC
));
if (sizeof_log > 1) PRINT(log);

But when I try to create my module I get an error that it is invalid.
If I add for example a raygen program to the same ptx the raygen program is identified but not the direct callables.

That’s most likely because the NVCC compiler command line is missing the option --relocatable-device-code=true (or short -rdc).
Compare with the setting I use in my examples: NVCC_OPTIONS

The reason why that is required since CUDA 8.0 is explained here:
https://forums.developer.nvidia.com/t/loading-callable-program-from-file-fails-with-rt-invalid-source/48648/2

Make sure to pull that OptiX 7 Course repository again. There has been a bug fix when using it with OptiX SDK 7.2.0 yesterday:
https://forums.developer.nvidia.com/t/optix-sdk-7-2-0-out-of-host-memory-error-7002/157558/3

Also note the other things I recommended to change in that function.
Use OPTIX_COMPILE_DEFAULT_MAX_REGISTER_COUNT.
Default-initialize all OptiX structures! You’re missing pipelineCompileOptions = {};
usesMotionBlur is an int type (The C API has no bool.)

1 Like

Thanks Detlef, that solved it! Always there when I need you.
Also thanks for your extra comments. I do initialise pipelineCompileOptions in the class declaration.

I’m now looking to use these callables in two separate pipelines, which preferably would have different pipelineCompileOptions as they don’t require the same amount of payload values. If I’ve understood correctly I’m required to use the same pipelineCompileOptions for all modules in a single pipeline as stated in the Programming Guide.

OptixPipelineCompileOptions
Must be identical for all modules used to create program groups linked in a single pipeline.

Am I correct in assuming that I’ll then have to create two different modules and subsequently two different program groups with the same callables if I want to use different pipelineCompileOptions in my pipelines? Worth mentioning that the ptx that hold the direct callables don’t hold any other programs so there is no optixTrace calls or anything in that specific module.

If you want to build different OptixPipeline objects, that would only be dependent on the OptixPipelineCompileOptions, the OptixPipelineLinkOptions and the list of OptixProgramGroup which reference the functions inside the OptixModules. See this optixPipelineCreate() call,

EDIT Oops, removed what I said before.
You’re correct. Since the OptixPipelineCompileOptions are used in optixModuleCreateFromPTX() as well, you would need to compile the modules twice if the OptixPipelineCompileOptions changed.
The debugLevel field should be matched in the OptixModuleCompileOptions and OptixPipelineLinkOptions.

Check this whole pipeline setup in my bigger example: Device::initPipeline()

1 Like

For people only reading the e-mails, ignore what I said about reusing OptixModules.

That won’t work when changing the OptixModuleCompileOptions or especially OptixPipelineCompileOptions for a different pipeline. See the edited text above.

I missed that when looking at the web sites. It’s easier to spot in a good text editor.