MDL in OptiX 7


I successfully converted some parts of my main application (using a PathTracer based on Detlefs Introduction Advanced Samples) to an OptiX7-based version. Denoiser, Callables, motionBlur, morphing, animation; all works great and seems really faster than before.
But I did not find any information about how to implement the OptiX direct callables (mdl_bsdf_*) of the MDL SDK in OptiX7. Its clear, where they need to be defined, but can they be simply accessed as any other callables?
And are functions as for example “mdl_helper->set_mdl_textures_ptx_string” still compatible?
Or is the MDL SDK version (MDL SDK 2018.1.2, build 312200.1281) which was shipped with OptiX 6.5 not even compatible at all with OptiX 7 ? Is there already a compatible MDL SDK version for OptiX 7 ?

Of course I wait for the next OptiX7 release (where AFAIK a MDL sample was promised), but it would be nice to know, whether some of the basics of that MDL interop has changed.

Thank you.


This sounds more like a question for the MDL SDK forum. Moving it there.

1 Like

Hi m1,

I didn’t check all versions, but you should be able to use MDL SDK 2019.1.1 or higher with OptiX 7.

With OptiX 7, we are not using callable programs for the texture runtime anymore, so you have to set the “tex_lookup_call_mode” backend option to “direct_call”. Thus “set_mdl_textures_ptx_string” is not used anymore.
As a starting point for the texture runtime, you can now use “texture_support_cuda.h” from MDL SDK examples/mdl_sdk/shared and include it in the CUDA file calling the functions generated by MDL SDK.
For loading the textures you can base on “Material_gpu_context::prepare_texture()” from “example_cuda_shared.h”.

This texture runtime uses the “res_data” parameter provided to the generated functions to access the textures. I put a pointer to the filled “Texture_handler” from “texture_support_cuda.h” in the hit group data.

mi::neuraylib::Resource_data res_data = {

So what about generating and calling the DF functions?
As you already expected, we are using direct callables for them:

We generate the target code and load the textures:

state.mdl_helper = new Mdl_helper(
    /*mdl_textures_ptx_path=*/ "",
    /*num_texture_spaces=*/ 1,
    /*num_texture_results=*/ 16 );

mi::base::Handle<mi::neuraylib::ITransaction> transaction =

std::vector<mi::neuraylib::Target_function_description> descs;

Material_info *mat_info = nullptr;
mi::base::Handle<const mi::neuraylib::ITarget_code> target_code(
        /*class_compilation=*/ true,

if (target_code->get_texture_count() > 1) {
    for (mi::Size i = 1, n = target_code->get_texture_count(); i < n; ++i) {

Create an OptiX module from the resulting PTX code:

OptixModuleCompileOptions moduleCompileOptions = {};
moduleCompileOptions.maxRegisterCount          = 0;
moduleCompileOptions.optLevel                  = OPTIX_COMPILE_OPTIMIZATION_DEFAULT;
moduleCompileOptions.debugLevel                = OPTIX_COMPILE_DEBUG_LEVEL_LINEINFO;

OptixPipelineCompileOptions pipelineCompileOptions = {};
pipelineCompileOptions.usesMotionBlur              = false;
pipelineCompileOptions.traversableGraphFlags       = OPTIX_TRAVERSABLE_GRAPH_FLAG_ALLOW_SINGLE_LEVEL_INSTANCING;
pipelineCompileOptions.numPayloadValues            = 2;
pipelineCompileOptions.numAttributeValues          = 2;
pipelineCompileOptions.exceptionFlags = OPTIX_EXCEPTION_FLAG_NONE;
pipelineCompileOptions.pipelineLaunchParamsVariableName = "params";

OptixModule module;
OPTIX_CHECK( optixModuleCreateFromPTX(
    &module ) );

Create a program group (per material) for the functions:

size_t callable_base_index = state.mdl_callables_groups.size();
state.mdl_callables_groups.resize(callable_base_index + 4);

OptixProgramGroupOptions callableProgramGroupOptions   = {};
OptixProgramGroupDesc    callableProgramGroupDesc[4]   = {};
callableProgramGroupDesc[0].kind                          = OPTIX_PROGRAM_GROUP_KIND_CALLABLES;
callableProgramGroupDesc[0].callables.moduleDC            = module;
callableProgramGroupDesc[0].callables.entryFunctionNameDC = "__direct_callable__mdlcode_init";
callableProgramGroupDesc[1].kind                          = OPTIX_PROGRAM_GROUP_KIND_CALLABLES;
callableProgramGroupDesc[1].callables.moduleDC            = module;
callableProgramGroupDesc[1].callables.entryFunctionNameDC = "__direct_callable__mdlcode_sample";
callableProgramGroupDesc[2].kind                          = OPTIX_PROGRAM_GROUP_KIND_CALLABLES;
callableProgramGroupDesc[2].callables.moduleDC            = module;
callableProgramGroupDesc[2].callables.entryFunctionNameDC = "__direct_callable__mdlcode_evaluate";
callableProgramGroupDesc[3].kind                          = OPTIX_PROGRAM_GROUP_KIND_CALLABLES;
callableProgramGroupDesc[3].callables.moduleDC            = module;
callableProgramGroupDesc[3].callables.entryFunctionNameDC = "__direct_callable__mdlcode_pdf";
OPTIX_CHECK( optixProgramGroupCreate(
    &state.mdl_callables_groups[callable_base_index] ) );

state.mdl_materials.push_back(MDLMaterial(mat_info, unsigned(callable_base_index)));


We save the “callable_base_index” here and later write it to the hit group data for the material together with the texture handler structure and the argument block of the material.

Add the program groups to the pipeline (when calling optixPipelineCreate) and add CallableSbtRecord objects for the to the SBT.

std::vector<CallableSbtRecord> callable_records;
for (OptixProgramGroup &group : state.mdl_callables_groups) {
    OPTIX_CHECK( optixSbtRecordPackHeader( group, &callable_records.back() ) );
CUdeviceptr d_callable_records = gpuMemDup(callable_records);


state.sbt.callablesRecordBase          = d_callable_records;
state.sbt.callablesRecordStrideInBytes = sizeof(CallableSbtRecord);
state.sbt.callablesRecordCount         = unsigned(callable_records.size());

In the renderer code, we then call the generated functions using “optixDirectCall” using IDs based on the callable_base_index stored in the hit groups:

const int callable_index_init = rt_data->mdl_callable_base_index;
const int callable_index_sample = rt_data->mdl_callable_base_index + 1;
const int callable_index_eval = rt_data->mdl_callable_base_index + 2;

optixDirectCall<void>(callable_index_init, &state, &res_data, /*exception_state=*/ nullptr, rt_data->arg_block);


optixDirectCall<void>(callable_index_eval, &eval_data, &state, &res_data, /*exception_state=*/ nullptr, rt_data->arg_block);


optixDirectCall<void>(callable_index_sample, &sample_data, &state, &res_data, /*exception_state=*/ nullptr, rt_data->arg_block);

I hope that helps as a starter and I hope I didn’t forget any relevant details.
If you need more information, just ask!

Best regards

1 Like

Hi Moritz,

thank you very much for your answer. I’ll go through that all.

This obviously requires a new version of mdl_helper.h . Where can I find that (until a new release of OptiX7 is present; the latest is from Aug 2019 and does not contain mdl_helper.h) ?
This file “mdl_helper.h” within the MDL SDK: is obviously not the one used here. And mdl_helper.h of OptiX 6.5 cannot be used for OptiX 7, cause it contains optix::Program and optix::Buffer objects. And there also was no function “compile_mdl_material()”. Instead “compile_expression()” was used to get the OptiX program.

So that means is not anymore used at all, right?
And so there is also no “texture_sampler_ids” buffer anymore, right?

When you refer to “the CUDA file”, do you mean a compiled CUDA kernel (linked as .OBJ) or the PTX file of the compiled OptiX-kernel?
AFAIK there was no CUDA kernel in OptiX 6.5 MDL handling. So I need an additional CUDA kernel for the texture access now?
So using “prepare_texture” will replace Mdl_helper::create_texture_samplers, right?
The texture samplers had a simple texture id; That one directly was accessible from an OptiX kernel. Is this still possible to access those textures from an OptiX7-kernel? The texture functions in texture_support_cuda.h are within the CUDA kernel (or are they part of the OptiX kernel?), can they be used from an OptiX7-kernel? I only need the albedo texture for the OptiX AI-denoiser and the alpha-component for a cutout-opacity anyhit program.
See your answer to my question from 2018: on 06/20/2018 :

Now that buffer of texture sampler ids is not there anymore. How to get access to them ?

In the OptiX 6.5 MDLSphere sample “res_data” was defined in device code; So now its instead in the hitgroup within the SBT. OK, I think I got that.

To be up-to-date I tried to get the latest MDL SDK 2019.2, cause says:
ABI compatible with the MDL SDK 2019.2 (325000.1814) binary release (see

But there is no MDL SDK 2019.2 (325000.1814) binary release on that page. Did I miss it?
only MDL SDK 2019.1.1 (317500.2554), so I’ll try that for now.

Thank you very much!


Hi m1,

Sorry for the confusion. As no OptiX 7 version of an MDL example has been released, yet, the new mdl_helper.h is not available, yet.
Let me summarize the differences between the pre-OptiX-7 version and current work-in-progress version of the mdl_helper:

  • The "tex_lookup_call_mode" option has been changed from "optix_cp" to "direct_call".
  • The texture handling has been changed from using optix::Buffer, optix::TextureSampler and rtTex2D/... to using cudaArray_t, cudaTextureObject_t and tex2D/.... Thus in mdl_helper the functions load_image_data(), load_texture(), create_texture_samplers() and set_texture_functions() are removed and replaced by copy_canvas_to_cuda_array() and prepare_texture() from example_cuda_shared.h.
  • No optix::Program and optix::Material objects are created anymore. Thus, all compile and create_program functions are removed. Instead the compile_mdl_material() function has been added, see below:
// Helper function to extract the module name from a fully-qualified material name.
std::string Mdl_helper::get_module_name(const std::string& material_name) const
    size_t p = material_name.rfind("::");
    return material_name.substr(0, p);

// Returns a string-representation of the given message severity
inline const char* message_severity_to_string(mi::base::Message_severity severity)
    switch (severity) {
    case mi::base::MESSAGE_SEVERITY_ERROR:
        return "error";
    case mi::base::MESSAGE_SEVERITY_WARNING:
        return "warning";
    case mi::base::MESSAGE_SEVERITY_INFO:
        return "info";
    case mi::base::MESSAGE_SEVERITY_VERBOSE:
        return "verbose";
    case mi::base::MESSAGE_SEVERITY_DEBUG:
        return "debug";
    return "";

// Returns a string-representation of the given message category
inline const char* message_kind_to_string(mi::neuraylib::IMessage::Kind message_kind)
    switch (message_kind) {
    case mi::neuraylib::IMessage::MSG_INTEGRATION:
        return "MDL SDK";
    case mi::neuraylib::IMessage::MSG_IMP_EXP:
        return "Importer/Exporter";
    case mi::neuraylib::IMessage::MSG_COMILER_BACKEND:
        return "Compiler Backend";
    case mi::neuraylib::IMessage::MSG_COMILER_CORE:
        return "Compiler Core";
    case mi::neuraylib::IMessage::MSG_COMPILER_ARCHIVE_TOOL:
        return "Compiler Archive Tool";
    case mi::neuraylib::IMessage::MSG_COMPILER_DAG:
        return "Compiler DAG generator";
    return "";

// Prints the messages of the given context.
// Returns true, if the context does not contain any error messages, false otherwise.
bool Mdl_helper::log_messages(mi::neuraylib::IMdl_execution_context* context)

    for (mi::Size i = 0; i < context->get_messages_count(); ++i) {
        mi::base::Handle<const mi::neuraylib::IMessage> message(context->get_message(i));
        m_last_mdl_error += message_kind_to_string(message->get_kind());
        m_last_mdl_error += " ";
        m_last_mdl_error += message_severity_to_string(message->get_severity());
        m_last_mdl_error += ": ";
        m_last_mdl_error += message->get_string();
        m_last_mdl_error += "\n";
    return context->get_error_messages_count() == 0;

// Compile the MDL material into target code.
mi::base::Handle<const mi::neuraylib::ITarget_code> Mdl_helper::compile_mdl_material(
    mi::neuraylib::ITransaction *transaction,
    std::string const &material_name,
    std::vector<mi::neuraylib::Target_function_description> &descs,
    bool class_compilation,
    Material_info **out_mat_info)
    mi::base::Handle<const mi::neuraylib::ITarget_code> code_ptx;

    // Create an execution context for options and error message handling
    mi::base::Handle<mi::neuraylib::IMdl_factory> mdl_factory(
    mi::base::Handle<mi::neuraylib::IMdl_execution_context> context(

    mi::base::Handle<mi::neuraylib::ILink_unit> link_unit(m_be_cuda_ptx->create_link_unit(
        transaction, context.get()));

    std::string module_name = get_module_name(material_name);
    m_mdl_compiler->load_module(transaction, module_name.c_str(), context.get());
    if (!log_messages(context.get()))
        return code_ptx;

    // Create a material instance from the material definition
    // with the default arguments.
    const char *prefix = (material_name.find("::") == 0) ? "mdl" : "mdl::";

    std::string material_db_name = prefix + material_name;
    mi::base::Handle<const mi::neuraylib::IMaterial_definition> material_definition(
    if (!material_definition) {
        m_last_mdl_error = "Material \"" + material_name + "\" not found";
        return code_ptx;

    mi::Sint32 ret = 0;
    mi::base::Handle<mi::neuraylib::IMaterial_instance> material_instance(
        material_definition->create_material_instance(0, &ret));
    if (ret != 0) {
        m_last_mdl_error = "Instantiating material \"" + material_name + "\" failed";
        return code_ptx;

    mi::Uint32 flags = class_compilation
        ? mi::neuraylib::IMaterial_instance::CLASS_COMPILATION
        : mi::neuraylib::IMaterial_instance::DEFAULT_OPTIONS;
    mi::base::Handle<mi::neuraylib::ICompiled_material> compiled_material(
        material_instance->create_compiled_material(flags, context.get()));
    if (!log_messages(context.get()))
        return code_ptx;

        compiled_material.get(),, descs.size(), context.get());
    if (!log_messages(context.get()))
        return code_ptx;

    code_ptx = mi::base::Handle<const mi::neuraylib::ITarget_code>(
        m_be_cuda_ptx->translate_link_unit(link_unit.get(), context.get()));
    if (!log_messages(context.get()))
        return code_ptx;

    if (out_mat_info != nullptr) {
        mi::base::Handle<mi::neuraylib::ITarget_value_layout const> arg_layout;
        mi::base::Handle<mi::neuraylib::ITarget_argument_block const> arg_block;

        if (code_ptx->get_argument_block_count() > 0) {
            arg_layout = code_ptx->get_argument_block_layout(0);
            arg_block = code_ptx->get_argument_block(0);

        *out_mat_info = new Material_info(

    return code_ptx;

The Material_info class is the Mdl_material_info class already released in examples\mdl_sdk\dxr\mdl_d3d12\mdl_material_info.h which helps editing material parameters.

Additionally, is replaced by texture_support_cuda.h which gets included in the .cu file which contains the shader calling the generated functions (for example a closesthit shader).

Well, I mean a .cu file which is used with optixModuleCreateFromPTX(). So it’s within the OptiX world and OptiX device functions are available.
To make it clear, if you would want to add MDL support to the optixPathTracer example, you would probably call the functions generated by the MDL SDK in the __closesthit__radiance in So in this file, I would include texture_support_cuda.h. You can probably also make the texture functions available in another way, though.

If you want, you can still access the textures using the same functions as in texture_support_cuda.h.
But since MDL SDK 2019.1.4 you can set the backend option “enable_auxiliary” to “on” to generate an additional “auxiliary” function which calculated the albedo and the normal for a hit point. Please refer to the df_cuda MDL SDK example to see how it works.

You do get the cutout opacity by generating and calling code for the MDL expression path “geometry.cutout_opacity”, don’t you?

In the MDLSphere example, res_data was a constant, so I made it a “device const”. Now, it’s a stack variable in the closesthit function, as it is not constant anymore but depends on the hitgroup.

The binary release of MDL SDK 2019.2 should be available on soon.

Best regards

1 Like

Hi Moritz,

thank you again very much for all the information. I think now I understand it.

Still I have one question about the: optixDirectCall-parameter “rt_data->arg_block”:
mdl_d3d12::Mdl_material_info in “mdl_material_info.h” clones mi::neuraylib::ITarget_argument_block const *arg_block to “m_arg_block”;
get_argument_block_data() returns a CPU pointer of the arg_block; get_argument_block_size() returns the size in bytes;
So all these blocks of all materials could be uploaded to one device buffer and so rt_data->arg_block
would be a device pointer into the ith material in that buffer (stored in SBT of that hitgroup), right? Or did I miss something?
is there any memory alignment required for those blocks on the GPU?

and one question about the “rt_data->texture_handler” in “res_data” parameter:
when calling “prepare_texture” then state.mdl_textures vector is expanded with “Texture” objects; those buffer contents need to be uploaded to GPU in a buffer. a device pointer to the first texture of a material then is used in “textures” field in “struct Texture_handler” in texture_support_cuda.h And so also that struct needs to be uploaded to the GPU. (so one Texture_handler struct and an offset into the textures-array per material on the GPU, right?) “rt_data->texture_handler” then is simply the device pointer into that ith struct Texture_handler of a material assigned in the hitgroup in the SBT, right?
When all the Texture_handler structs are applied in one device buffer, does each of then need to be aligned on a memory boundary ?
In your code there is also a vector of materials in the state:

state.mdl_materials.push_back(MDLMaterial(mat_info, unsigned(callable_base_index)));

“MDLMaterial” seems to be a custom struct, right?
Each material then would have its own Texture_handler referencing its “Texture” objects, right?

No, in my OptiX 6.5-based implementation, my anyhit shader looks similar to this one:,
but with the difference, that for opacity only the alpha channel from a float4 diffuse texture (which also contains a float3 albedo color) is used (accessed through texture id .getId() from texture sampler).
That cutout opacity is not defined in the MDL material. If that direct texture access would not be possible this way, I would need to load that texture again. However, now when its defined in the Texture_sampler I only need the index of it in that “Texture” vector. That index I think I will get as before by traversing through all target_code->get_texture(i) file names. The albedo filename then is the desired file. Using that texture_idx parameter on calling tex_texel_float4_2d will give me the float4 color, right?

This is great news! Will try that when the binary release of MDL SDK 2019.2 is available.

thank you.

Yes, you can do that. We’d recommend 16 byte alignment per argument block, should be enough for now.

Actually, the textures-array is per link-unit, so if you have only one material per link-unit, yes, you have one textures-array per material ;)
But yes, you could also put them all into one buffer.

The generated code actually does not access the Texture_handler structs if the “tex_lookup_call_mode” is not “vtable”, it is just provided to the texture handler functions untouched. So the required alignment only depends on your texture handler implementation.

Yes and depends on the number of materials per link-unit.

Yes, this would work (but this gives you the unfiltered result, use “tex_lookup_float4_2d” for the filtered one).
Sounds like quite a workaround. Were there any problems defining the cutout opacity in the MDL material (like “cutout_opacity: tex::lookup_float4(albedo_texture, coords).w”) and generating code to evaluate this expression?

Best regards

1 Like

The anyhit shader only needs one direct texture lookup this way.
And so no additional code generated from the evaluated expression of the MDL material is needed. And also no additional call to a “direct callable” would be required.
I think its faster when not calling such an additional “direct callable”. And yet I’m also not aware of a MDL example (for OptiX), which could operate on OptiX-geometry using “geometry.cutout_opacity”. What type of direct callables would be provided in that case?
Would it also be “__direct_callable__mdlcode_init”? And how should that interact then with an OptiX anyhit shader? Is there even a way to apply it in an OptiX-Intersection shader?

thank you,

Yes, you’re right, there seems to be some overhead you can avoid this way. Also you won’t need to initialize the pretty big MDL state structure.

With “geometry.cutout_opacity” I’m refering to the MDL material “cutout_opacity” field in “material_geometry” (see section 13.3 in the MDL 1.6.1 Spec), not to an OptiX entity.
If you add another “Target_function_description”


MDL SDK will generate the function “__direct_callable__mdlcode_cutout_opacity” with the type “mi::neuraylib::Material_expr_function” defined in “mi/neuraylib/target_code_types.h”.
If you have the index of the function “callable_index_cutout”, you can call it like that in a shader:

float cutout_opacity = 0.0f;
optixDirectCall<void>(callable_index_cutout, &cutout_opacity, &state, &res_data, nullptr, rt_data->arg_block);

Setting up the MDL state might be more cumbersome, but I don’t see a general problem here. At least according to, it is possible to call direct callables from intersection programs. Maybe I’m missing something, though.

1 Like

Thank you for the information.

Before implementing OptiX7-handling I first rebuilt my OptiX6.5-based project from MDL SDK 2018.1.2, build 312200.1281 to MDL SDK 2019.1.1, build 317500.2554; This works withoug problems.
The handling in my module yet uses basically the technique as shown in the optixMDLSphere sample of OptiX 6.5.0, with some additional handling in the device code (based on data.sample.event_type contents);

But today I rebuilt that project again with the MDL SDK 2019.2, build 325000.1814 binary release (and also updating compiled_material_traverser_base.h/.cpp and compiled_material_traverser_print.h/.cpp).

For test I used in both cases gun_metal.mdl (in\mdl-sdk-325000.1814\examples\mdl\nvidia\sdk_examples)
But the output seems invalid for the newer version (see attachment REMOVED).
I’ve not found any documented changes related to the OptiX-6.5 interop.
So for now I think I will try to implement OptiX 7 using MDL SDK 2019.1.1, build 317500.2554.

NOTE: the new versions of compiled_material_traverser_* work also properly with the MDL SDK 2019.1.1

My system: OptiX 6.5.0 SDK CUDA 10.1.243 GTX 1050 2GB Win10PRO 64bit (version 1809; build 17763.107) device driver: 436.48 VS2019 v16.4.2 (toolkit v140 of VS2015)

Obviously the new MDL SDK sends some other expression states (data.sample.event_type & mi::neuraylib::BSDF_EVENT_*) on the device code. Is there a list of such changes between 2019.1.1 and 2019.2 ?
Has especially “simple_glossy_bsdf” 's event firing changed?

Or maybe its a bug in my code, which was simply not visible in the old version.

I quickly ported the OptiX 6.5 optixMDLSphere example to MDL SDK 2019.2 and I only applied these changes to

We don’t want to use light path expressions (and we didn’t enable them in the backend via the “df_handle_slot_mode” option):

mi::neuraylib::Bsdf_evaluate_data evaluate;


mi::neuraylib::Bsdf_evaluate_data<mi::neuraylib::DF_HSM_NONE> evaluate;

The multiscatter_tint support requires a forth random value for sampling:

data.sample.xi = make_float3(xi0, xi1, xi2);

mdl_bsdf_sample(&data.sample, &state, &res_data, NULL, NULL);

const float xi3 = rnd(prd_radiance.rnd_seed);

data.sample.xi = make_float4(xi0, xi1, xi2, xi3);

mdl_bsdf_sample(&data.sample, &state, &res_data, NULL, NULL);

mi::neuraylib::Bsdf_evaluate_data now always splits between diffuse and glossy contibution, so add them up:

prd_radiance.result += f * data.evaluate.bsdf * mis_weight;


prd_radiance.result += f * (data.evaluate.bsdf_diffuse + data.evaluate.bsdf_glossy) * mis_weight;

The result looks correct (see attachment). Even multiscatter_tint works with the MDL helper (actually just by chance ^_^ mi::neuraylib::ITarget_code::Texture_shape_bsdf_data should be handled just like mi::neuraylib::ITarget_code::Texture_shape_3d).

I compared the mi::neuraylib::Bsdf_event_type enum from 2018.1 and 2019.2 and there are no differences. Which additional handling do you mean?

1 Like

Oh! Thank you very much for that update.
I was not aware of these changes. I simply used the already compiled OptiX-kernels with the new MDL SDK, which obviously is not possible. Now its clear why. After re-compilation on the OptiX6.5-based version now everything works fine with the new MDL SDK 2019.2 .

You’re very welcome!
If you have any further questions regarding OptiX 7 and MDL SDK 2019.2, just ask!

One important detail which cost me quite some time debugging:
Don’t forget to adapt the maxDCDepth parameter of optixUtilComputeStackSizes() for the additional direct callable calls and also call optixUtilAccumulateStackSizes() for the MDL callables group. Otherwise you may get hangs or wrong rendering results.

1 Like

Hi Moritz,

thanks for the info.

I tried to implement it now. In yet for test I simply exit without any calls; pure solid color.

On running a test using ::example_df::df_material I get this error:

[3324] Optix call optixPipelineCreate( context, ADDRpipeline_compile_options, &pipeline_link_options, program_groups, numProgramGroups, log, &sizeof_log, ADDRpipeline ) failed: line:4348 Error Code: 7251
[3324] Log: COMPILE ERROR: failed to create pipeline
[3324] Info: Pipeline has 5 module(s), 9 entry function(s), 1 trace call(s), 0 continuation callable call(s), 0 direct callable call(s), 221 basic block(s) in entry functions, 5803 instruction(s) in entry functions, 0 non-entry function(s), 0 basic block(s) in non-entry functions, 0 instruction(s) in non-entry functions


When the MDL-material DC program groups were not added to program_groups, pipeline creation succeeded.

It also succeeds with the MDL-DC’s, when running a test on the same code with ::gun_metal::gun_metal or ::carpaint_orange_peel::carpaint_orange_peel_3_layer

The only difference I see there is, that ::example_df::df_material uses a texture, but the other 2 don’t.

So what could be wrong in my code with the pipeline linking, when there is a texture in the MDL ?
In both cases exactly the same code is used, so the MDL callables were only different.

thank you.


I love such helpful error messages…

Good observations you made there already! From this my guess would be, that the linker cannot find the texture runtime functions.

Did you do this? I have the include in the same .cu file which contains the closesthit program doing the optixDirectCall invocations. So in your case it should probably be included in

1 Like

indeed, I thought I had put it there, but it was in another .cu; so it was missing; thanks!
I also enabled the mdl handling (removing the solid color test output).

But now at run-time the compiled PTX of the file fails with this error:

[11164] Optix call optixModuleCreateFromPTX( state.context, &module_compile_options, &state.pipeline_compile_options, ptx_string, ptx_size, log, &sizeof_log, ptx_module ) failed:
line:3796 Error Code: 7200
[11164] Log: COMPILE ERROR: Invalid PTX input: ptx2llvm-module-003: error: Failed to translate PTX input to LLVM
[11164] PTX symbols of certain types (e.g. pointers to functions) cannot be used to initialize array “tex_vtable”.
[11164] PTX symbols of certain types (e.g. pointers to functions) cannot be used to initialize array “tex_deriv_vtable”.
get symbol failed: __cudart_i2opi_f
get symbol failed: __cudart_i2opi_f
get symbol failed: __cudart_i2opi_f
get symbol failed: __cudart_i2opi_f
get symbol failed: __cudart_i2opi_f
get symbol failed: __cudart_i2opi_f
[11164] get symbol failed: __cudart_i2opi_f
[11164] get symbol failed: __cudart_i2opi_f
[11164] get symbol failed: __nv_static_54__43_tmpxft_000013b8_00000000_10_mdl_ibl_cpp1_ii_params_identity
[11164] get symbol failed: __cudart_i2opi_f
[11164] get symbol failed: __cudart_i2opi_f
[11164] get symbol failed: __cudart_i2opi_f
[11164] get symbol failed: __cudart_i2opi_f


My system: OptiX 7.0.0 SDK CUDA 10.1.243 GTX 1050 2GB Win10PRO 64bit (version 1809; build 17763.107) device driver: 436.48 VS2019 v16.4.2 (toolkit v140 of VS2015) was compiled this way (and succeeded):

@echo "%CUDA_PATH%"
@set PATH=%PATH%;"C:\Program Files (x86)\Microsoft Visual Studio 14.0\VC\bin"

"%CUDA_PATH%\bin\nvcc.exe"  -ptx   -ftz true -O2  --cudart static --machine 64 -arch compute_60 --relocatable-device-code=true

(I generally do not use fast math for precision)

I also for test copied cudart64_101.dll into the same folder, where the exe file resides.
No change.

Thank you.

Sorry, I thought I mentioned, that you need to comment out the vtables at the end of “texture_support_cuda.h”, but looks like I didn’t. This is causing the “PTX symbols of certain types” errors.

I cannot reproduce the “__cudart_i2opi_f” errors, though.
But this failed symbol lookup also looks weird “__nv_static_54__43_tmpxft_000013b8_00000000_10_mdl_ibl_cpp1_ii_params_identity”, as the name implies that this symbol should maybe be defined in your code.
Maybe it’s a follow-up error?

If not, maybe my nvcc commandline contains any relevant difference (generated by the adapted CMake files of the OptiX SDK 7.0.0 examples):

"C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v10.1/bin/nvcc.exe" "C:/ProgramData/NVIDIA Corporation/OptiX SDK 7.0.0/SDK/optixMDL/" -ptx -o D:/source2/optix-7.0.0-build/lib/ptx/ -ccbin "C:/Program Files (x86)/Microsoft Visual Studio/2017/Professional/VC/Tools/MSVC/14.16.27023/bin/Hostx64/x64" -m64 -D_USE_MATH_DEFINES -DNOMINMAX -arch sm_30 --use_fast_math --compiler-options /D_USE_MATH_DEFINES -rdc true -DNVCC "-IC:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v10.1/include" "-IC:/ProgramData/NVIDIA Corporation/OptiX SDK 7.0.0/include" "-IC:/ProgramData/NVIDIA Corporation/OptiX SDK 7.0.0/SDK/cuda" "-IC:/ProgramData/NVIDIA Corporation/OptiX SDK 7.0.0/SDK" -ID:/source2/optix-7.0.0-build/include "-IC:/ProgramData/NVIDIA Corporation/OptiX SDK 7.0.0/SDK/support/mdl-sdk/include" -ID:/source2/optix-7.0.0-build "-IC:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v10.1/include"
1 Like

Thanks ! Commenting out that part did the trick. The other one obviously was a follow-up error, cause now module and pipeline creation succeeds. And also rendering succeeds. (again only yet using solid color in the closesthit for test and not actually calling the MDL-related DC’s; Will try that later on.)

I finally got it done now. MDL materials run under OptiX 7.0.0 in my app.
Great Thanks to you, Moritz!!!