MDL malformed PTX input

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

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.

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

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).