OptiX 7 and PTX interop

Gaming and Visualization Technologies Visualization OptiX
1

Hi,

I have some external PTX modules, I’d like to use them as direct callables, so that these external work can be used in OptiX. By using optixModuleCreateFromPTX and entryFunctionNameDC, I can register functions to the pipeline, that’s great.

However, I don’t know how to assign the variables of these PTX modules. With cuda, I can use cuModuleGetGlobal. Unfortunately, there seems no way to get cuModule from OptiXModule.

As OptiX 7 is very cuda like, are there any way to assign the variables of external PTX modules?

Example:

.visible .global .align 4 .b8 diffuseColor[12];
.visible .global .align 4 .f32 diffuseRoughness;
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo12diffuseColorE[8] = {82, 97, 121, 0, 12, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo16diffuseRoughnessE[8] = {82, 97, 121, 0, 4, 0, 0, 0};

	// .globl	__direct_callable__aaa
.visible .func  (.param .align 4 .b8 func_retval0[12]) __direct_callable__aaa(
	.param .b64 __direct_callable__aaa_param_0,
	.param .b64 __direct_callable__aaa_param_1
)
{
	.reg .f32 	%f<4>;

	ld.global.f32 	%f1, [diffuseColor];
	ld.global.f32 	%f2, [diffuseColor+4];
	ld.global.f32 	%f3, [diffuseColor+8];
	st.param.f32	[func_retval0+0], %f1;
	st.param.f32	[func_retval0+4], %f2;
	st.param.f32	[func_retval0+8], %f3;
	ret;
}

	// .globl	__direct_callable__bbb
.visible .func  (.param .align 4 .b8 func_retval0[12]) __direct_callable__bbb(
	.param .b64 __direct_callable__bbb_param_0,
	.param .b64 __direct_callable__bbb_param_1
)
{
	.reg .f32 	%f<2>;

	ld.global.f32 	%f1, [diffuseRoughness];
	st.param.f32	[func_retval0+0], %f1;
	st.param.f32	[func_retval0+4], %f1;
	st.param.f32	[func_retval0+8], %f1;
	ret;
}

In this case, I don’t know how to set diffuseColor and diffuseRoughness.

Thanks

1 Like
2

The ideal thing to do is put those global variables inside your launch parameters in constant memory instead. Do you have anything preventing that? You would declare them in your module’s code the same way they’re declared in any OptiX program.

–
David.

3

Hi David,

Thank you for the answer.

It would be tricky for me to convert those global variables to constant memory, because those PTX modules are generated by third parties. In fact the function name is also a problem, I changed that by replacing the name with prefix direct_callable. The example is a simplest case for the remaining problem (the global variable).

Are there any non-ideal approaches to tackle this issue?

4

Ah, I see, that’s a tough case. Are you able to negotiate any changes or interface to the 3rd party PTX, or is this code you can’t realistically talk to the authors about at all?

There isn’t a way to register and write to variables on the host side for OptiX modules, so here is one possible option that I speculate might work, with the caveat that I haven’t tried it. I hope it’s helpful and that I’m not sending you down a rabbit hole.

Probably the easiest of the very ugly solutions would be to setup a separate “setter” launch of size 1 to write all your globals from host memory into device memory, before you run the “render” launch.

The setter launch would need it’s own separate raygen program, with the global vars linked as extern, that does all the copying of data from launch params to your external 3rd party global vars. (Or, if you’d rather, you can use your own cuda buffer instead of launch params.)

You’ll need a separate SBT for the setter launch that runs your setter raygen program.

For it all to work, the setter code and the render code needs to be compiled together in the same pipeline. So, you want 1 pipeline, 2 SBTs, 2 raygen programs, and 2 sets of launch params.

I am imagining that if this works, you could abstract it into a host side function that takes a list of global var names and sizes, and then writes & compiles the setter raygen program on the fly, and then runs the synchronous setter launch, all in one go.

–
David.

5

Incidentally, I heard that function parameters do not require the direct_callable tag, only the function names do. Our examples do that by convention mostly for clarity, but maybe it’ll save you some work if you don’t have to rename the parameters.

–
David.

6

It works. Thank you David.