MDL malformed PTX input

Gaming and Visualization Technologies Visualization MDL SDK
1

I’m looking to compile an MDL file and use it within Optix7. I compiled to PTX and when trying to create an Optix7 module, I’m getting this message:

OPTIX COMPILE FEEDBACK => COMPILE ERROR: No functions with semantic types found

Malformed PTX input. See compile details for more information.
Info: Module uses 0 payload values. Pipeline configuration: 6.
Info: Module uses 0 attribute values. Pipeline configuration: 3.
Info: 11 non-entry function(s) have 67 basic block(s), 1115 instruction(s)

Plus, when looking at the PTX code, the direct call functions are defined as

.visible .func __direct_call__X_init

But when I look into the PTX file of a *.cu file directly compiled with nvrtcCompileProgram, I’m getting functions defined as

.visible .entry __closesthit__red()

which is a kernel function.

Not sure what I’m missing here.

Jack

2

I found some information in another blog post (I didn’t kept the link sadly) that suggested to add something like:

.visible .entry __direct_callable__dummy__some_unique_suffix() {ret;}

And now it compiles and I can load it.

The only problem now, is after calling _init, _sample, and _eval, in that order, I cannot read the simple color value that my MDL should return.
The sample steps also doesn’t alter the value of the event_type, and some problem with _eval that doesn’t change the bsdf_diffuse.

3

A little bit more information is needed. Does this happen with the official MDL example that comes with OptiX?

4

I’m currently trying to integrate MDL into an Optix7 project I made from scratch (mostly learning various aspect of RTX).

I started with a very simple material, with the objective of just setting the base color:

export material testing(
uniform color baseColor= color(1.0, 0.75, 0.0)
) = let
{
bsdf diff = df::diffuse_reflection_bsdf(tint : baseColor);
} in material(
surface: material_surface(
scattering: diff
)
);

I solved the call of the direct callables that I had, but now it’s a memory error.

This is with cuda-gdb:

[Switching focus to CUDA kernel 0, grid 43, block (2751,0,0), thread (0,0,0), device 0, sm 0, warp 2, lane 0]
0x0000555559369368 in __direct_callable__X_sample_0x3cf3d8de778f880f<<<(5120,3,1),(64,1,1)>>> ()

And cuda-memcheck:

Out-of-range Shared or Local Address
at 0x00000030 in __direct_callable__X_sample_0x3cf3d8de778f880f
by thread (31,0,0) in block (2732,0,0)

And no idea how to check which memory access is wrong.

This is part of what I have so far:

mi::neuraylib::Shading_state_material state = {
normal,
normal,
normal, // the object is a unit sphere
0.0f,
&uv,
&tangent,
&binormal,
nullptr, // texture result
nullptr, // RO data segment
identity, // world to object
identity, // object to world
0, // object ID
1.0f // meter per unit
};

// texHandler is a nullptr for the vtable and compiled in direct_call mode
mi::neuraylib::Resource_data resData = { nullptr, &texHandler };

mdlcode_init(state, &resData, nullptr, args);

mi::neuraylib::Bsdf_sample_data sampleData;
sampleData.ior1 = make_float3(1.0f);
sampleData.ior2.x = MI_NEURAYLIB_BSDF_USE_MATERIAL_IOR;
sampleData.k1 = -worldDir;
sampleData.xi = make_float4(0.0f);

mdlcode_sample(&sampleData, state, &resData, nullptr, args);

And then it crashes.

As for using more complex materials, that will be later work. Need to extract the arguments (right now it’s hardcoded), add texture support, extract the constants values (like thin walled when there’s no callables for it).

5

From this code snipped, I cannot see any errors. Ensure that the constant data segment generation is disabled if you pass nullptr here.

Do you compile in instance compilation mode or in class-compilation mode? Try instance compilation first, as this further reduces dependencies on memory access.

6

I added the compile_constants option to off just like you suggested and tried the class compilation mode and still no changes.

But in general, it still points to an memory access error.

But now, with cuda-gdb, I’m getting a new error:

in state::normal()<<<(5120,3,1),(64,1,1)>>> ()

even if the normal is not altered yet.

7

Hi Jack,

I noticed one problem, which may result in a crash depending on the backend options you are using: you’re providing a nullptr for texture results in your Shading_state_material struct.
Usually, the init function (mdlcode_init in your case) will precalculate some expressions (not only textures, but also more expensive calculations), which will be used by the sample, evaluate and pdf functions. You can set the size of the texture results in number of float4s with the IMdl_backend option “num_texture_results”. The default is zero, but you probably have set this to something else for better performance. The OptiX7 example sets the option to 16.
In the CUDA code, it uses a local variable for the texture results array:

float4 texture_results[16];
Mdl_state state;
[...]
state.text_results = texture_results;

Does this solve your crash problem?

For the initial problem, you used the wrong semantic prefix for the functions, right? The prefix should be

"__direct_callable__"

Then at least I don’t need any dummy function.

Please also note, that not using direct callables for calling the MDL code may significantly improve performance. The optix7 example of the MDL SDK uses the “llvm_renderer_module” option to optimize the generated code together with the renderer part calling the MDL code. This example also can be configured to use direct callables for comparison.
The example also contains a GUI for changing material arguments at runtime.

Hope that helps
Moritz