MDL SDK API nvidia_logo_transpbg.gif Up
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Groups Pages
Texture lookup call modes of the PTX backend

The PTX backend allows to configure how texture lookup methods are called in the generated code.

The mode can be chosen by setting the "tex_lookup_call_mode" option via the mi::neuraylib::IMdl_backend::set_option() method to one of the modes described in the next sections:

  • "vtable"
  • "direct_call"
  • "optix_cp"

Depending on the mode, the generated code makes different use of mi::neuraylib::Resource_data parameter res_data which has to be provided to the generated code by the renderer:

struct Resource_data {
void const *shared_data;
const Texture_handler_base *texture_handler;
};
Note
The shared_data field is currently unused and should always be NULL.

"vtable" mode

This is the default mode. In this mode, the texture_handler field of the res_data parameter has to point to a customizable structure that can be derived from mi::neuraylib::Texture_handler_base.

The only requirement is that the first field of this structure is a pointer to a vtable (table of function pointers) pointing to the texture access functions (see below). Any additional data needed by your implementation can be appended after the vtable. The texture access functions can access the structure via the self parameter.

The retrieved texture value should be stored in the result parameter.

The texture_idx parameter is the index of the texture resource according to the mi::neuraylib::ITarget_code::get_texture() function. Note that index 0 is always the invalid texture.

Please refer to the MDL specification (section 20.3) for an explanation of the other parameters.

// The wrap mode determines the texture lookup behavior if a lookup coordinate
// is exceeding the normalized half-open texture space range of [0, 1).
enum Tex_wrap_mode {
wrap_clamp = 0, // clamps the lookup coordinate to the range
wrap_repeat = 1, // takes the fractional part of the lookup coordinate
// effectively repeating the texture along this axis
wrap_mirrored_repeat = 2, // like wrap_repeat but takes one minus the fractional
// part every other interval to mirror every second
// instance of the texture
wrap_clip = 3 // makes the texture lookup return zero for texture
// coordinates outside of the range
};
// The vtable for the texture handler.
void (*m_tex_lookup_float4_2d)(
float result[4],
Texture_handler_base const *self,
unsigned texture_idx,
float const coord[2],
Tex_wrap_mode wrap_u,
Tex_wrap_mode wrap_v,
float const crop_u[2],
float const crop_v[2]);
void (*m_tex_lookup_float3_2d)(
float result[3],
Texture_handler_base const *self,
unsigned texture_idx,
float const coord[2],
Tex_wrap_mode wrap_u,
Tex_wrap_mode wrap_v,
float const crop_u[2],
float const crop_v[2]);
void (*m_tex_texel_float4_2d)(
float result[4],
Texture_handler_base const *self,
unsigned texture_idx,
int const coord[2],
int const uv_tile[2]);
void (*m_tex_lookup_float4_3d)(
float result[4],
Texture_handler_base const *self,
unsigned texture_idx,
float const coord[3],
Tex_wrap_mode wrap_u,
Tex_wrap_mode wrap_v,
Tex_wrap_mode wrap_w,
float const crop_u[2],
float const crop_v[2],
float const crop_w[2]);
void (*m_tex_lookup_float3_3d)(
float result[3],
Texture_handler_base const *self,
unsigned texture_idx,
float const coord[3],
Tex_wrap_mode wrap_u,
Tex_wrap_mode wrap_v,
Tex_wrap_mode wrap_w,
float const crop_u[2],
float const crop_v[2],
float const crop_w[2]);
void (*m_tex_texel_float4_3d)(
float result[4],
Texture_handler_base const *self,
unsigned texture_idx,
int const coord[3]);
void (*m_tex_lookup_float4_cube)(
float result[4],
Texture_handler_base const *self,
unsigned texture_idx,
float const coord[3]);
void (*m_tex_lookup_float3_cube)(
float result[3],
Texture_handler_base const *self,
unsigned texture_idx,
float const coord[3]);
void (*m_tex_resolution_2d)(
int result[2],
Texture_handler_base const *self,
unsigned texture_idx,
int const uv_tile[2]);
};
// Your texture handler structure.
struct Texture_handler : public Texture_handler_base {
// additional data for the texture access functions can be provided here
int example_field;
};

"direct_call" mode

In this mode, the texture access functions will be called directly in the PTX code. The implementations must be linked to the generated code.

The expected prototypes correspond to the prototypes in the Texture_handler_vtable structure of the vtable, except that the Texture_handler structure is not required to contain the vtable field. The texture_handler field of the mi::neuraylib::Resource_data structure may be NULL, if your implementation does not need it.

Here are the CUDA prototypes of the functions you have to implement:

// Your texture handler structure.
struct Texture_handler : public Texture_handler_base {
// additional data for the texture access functions can be provided here
int example_field;
};
extern "C" __device__ void tex_lookup_float4_2d(
float result[4],
Texture_handler const *self,
unsigned texture_idx,
float const coord[2],
Tex_wrap_mode wrap_u,
Tex_wrap_mode wrap_v,
float const crop_u[2],
float const crop_v[2]);
extern "C" __device__ void tex_lookup_float3_2d(
float result[3],
Texture_handler const *self,
unsigned texture_idx,
float const coord[2],
Tex_wrap_mode wrap_u,
Tex_wrap_mode wrap_v,
float const crop_u[2],
float const crop_v[2]);
extern "C" __device__ void tex_texel_float4_2d(
float result[4],
Texture_handler const *self,
unsigned texture_idx,
int const coord[2]);
extern "C" __device__ void tex_lookup_float4_3d(
float result[4],
Texture_handler const *self,
unsigned texture_idx,
float const coord[3],
Tex_wrap_mode wrap_u,
Tex_wrap_mode wrap_v,
Tex_wrap_mode wrap_w,
float const crop_u[2],
float const crop_v[2],
float const crop_w[2]);
extern "C" __device__ void tex_lookup_float3_3d(
float result[3],
Texture_handler const *self,
unsigned texture_idx,
float const coord[3],
Tex_wrap_mode wrap_u,
Tex_wrap_mode wrap_v,
Tex_wrap_mode wrap_w,
float const crop_u[2],
float const crop_v[2],
float const crop_w[2]);
extern "C" __device__ void tex_texel_float4_3d(
float result[4],
Texture_handler const *self,
unsigned texture_idx,
int const coord[3]);
extern "C" __device__ void tex_lookup_float4_cube(
float result[4],
Texture_handler const *self,
unsigned texture_idx,
float const coord[3]);
extern "C" __device__ void tex_lookup_float3_cube(
float result[3],
Texture_handler const *self,
unsigned texture_idx,
float const coord[3]);
extern "C" __device__ void tex_resolution_2d(
int result[2],
Texture_handler const *self,
unsigned texture_idx,
int const uv_tile[2]);
Note
You have to declare the implementations as extern "C" when using CUDA code to avoid name mangling.

"optix_cp" mode

This mode is meant for applications using the OptiX 6 framework (for OptiX 7, the "direct_call" mode is recommended). In this mode, the generated code will contain OptiX rtCallableProgramId variables which have to be set to the texture access functions like this (mdl_textures_ptx_path is the path to a PTX file containing the implementations of the texture access functions, mdl_expr_prog is the OptiX program created from the PTX code generated for the material sub-expression):

// Set up the texture access functions for the MDL expression code.
static const char *tex_prog_names[] = {
"tex_lookup_float4_2d",
"tex_lookup_float3_2d",
"tex_texel_float4_2d",
"tex_lookup_float4_3d",
"tex_lookup_float3_3d",
"tex_texel_float4_3d"
"tex_lookup_float4_cube",
"tex_lookup_float3_cube",
"tex_resolution_2d"
};
for (size_t i = 0; i < sizeof(tex_prog_names) / sizeof(*tex_prog_names); ++i) {
optix::Program tex_prog = m_optix_ctx->createProgramFromPTXFile(
mdl_textures_ptx_path, tex_prog_names[i]);
// your code should set any required variables in the texture access program here,
// like a buffer of texture sampler IDs for the used textures
tex_prog["texture_sampler_ids"]->set(tex_samplers);
mdl_expr_prog[tex_prog_names[i]]->setProgramId(tex_prog);
}

In this example, the names of the methods in the PTX file referenced by mdl_textures_ptx_path are identical to the names of the needed OptiX variables, but this is not required.

The prototypes of the functions are identical to the "direct_call" mode except that extern "C" __device__ is replaced by RT_CALLABLE_PROGRAM.