Material Definition Language API nvidia_logo_transpbg.gif Up
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Groups Pages
Example for Texture Filtering with Automatic Derivatives
[Previous] [Up] [Next]

This example shows how to enable and use automatic derivatives for texture filtering.

New Topics

  • Automatic derivatives

Detailed Description

Automatic derivatives


When rendering images with high-frequency content, like a checker board texture, you usually have to shoot an increased number of rays to avoid aliasing effects. By providing information to the texture lookup functions about the area covered by a ray, the integral over this area can be calculated (or approximated) to get good texture sampling with just a single ray.

In the MDL API, the renderer can provide the derivatives of texture coordinates with respect to x and y in screen space, when calling generated code for MDL expressions or distribution functions. The generated code then calculates the derivatives of expressions provided to 2D texture lookup functions as texture coordinate parameters using "Automatic Differentiation" (see Dan Piponi, "Automatic Differentiation, C++ Templates, and Photogrammetry", Journal of Graphics, GPU, and Game Tools (JGT), Vol 9, No 4, pp. 41-55 (2004)). The renderer-provided texture runtime can finally use texture filtering techniques like elliptically weighted average (EWA) to determine the texture result.

Automatic derivatives are supported by the CUDA, LLVM IR and native backends of the MDL API. To enable this feature, you need to:

  • set the backend option "texture_runtime_with_derivs" to "on" via the IMdl_backend::set_option() method
  • replace Shading_state_material by Shading_state_material_with_derivs and provide the derivatives of the texture coordinates
  • provide a texture runtime supporting "tex_lookup_deriv_float4_2d" and "tex_lookup_deriv_float3_2d" (for the native backend, there is a simple isotropic implementation using mipmaps, if you use the default "on" value for the "use_builtin_resource_handler" option)

In the "example_df_cuda" example, the backend option is set in the constructor of the Material_compiler class in example_cuda_shared.h, when the "-d" option is provided via command line:

check_success(m_be_cuda_ptx->set_option("texture_runtime_with_derivs", "on") == 0);

In this mode, the example will use a variant of example_df_cuda.cu compiled with a "ENABLE_DERIVATIVES" define, which will enable calculating the texture coordinate derivatives on the sphere. It does so by intersecting two additional rays, offset by one-pixel in x and y direction, with the plane given by the primary intersection point and the corresponding normal. Using the surface derivatives with respect to U and V and the new intersection points, it determines the derivatives of the texture coordinates with respect to screen-space x and y. See Matt Pharr et al., "Physically Based Rendering", 3rd edition (2016), chapter 10.1.1 for details.

Instead of the usual Shading_state_material, the example then fills in the derivative variant Shading_state_material_with_derivs where the texture coordinates are tct_deriv_float3 values consisting of val, the texture coordinate, dx and dy, the derivative of the texture coordinate with respect to screen-space x and screen-space y, respectively.

The types could be written as:

tct_float3 val; // value component
tct_float3 dx; // derivative of value with respect to screen-space x
tct_float3 dy; // derivative of value with respect to screen-space y
};
tct_float3 normal; // state::normal() result
tct_float3 geom_normal; // state::geom_normal() result
tct_float3 position; // state::position() result
tct_float animation_time; // state::animation_time() result
const tct_deriv_float3 *text_coords; // state::texture_coordinate() table
const tct_float3 *tangent_u; // state::texture_tangent_u() table
const tct_float3 *tangent_v; // state::texture_tangent_v() table
tct_float4 *text_results; // texture results lookup table
const char *ro_data_segment; // read-only data segment
const tct_float4 *world_to_object; // world-to-object transform matrix
const tct_float4 *object_to_world; // object-to-world transform matrix
tct_int object_id; // state::object_id() result
};
Note
For the native backend, you have to cast the Shading_state_material_with_derivs object to a Shading_state_material reference when calling the "execute_*" methods of mi::neuraylib::ITarget_code, as these methods have not been duplicated for the derivative variants.

Additionally, a different vtable "tex_deriv_vtable" defined in texture_support_cuda.h is provided for the texture handler, which contains derivative variants for the 2D texture lookup functions: The functions "tex_lookup_deriv_float4_2d" and "tex_lookup_deriv_float3_2d" expect a "tct_deriv_float2 const *" as "coord" parameter:

tct_float2 val; // value component
tct_float2 dx; // derivative of value with respect to screen-space x
tct_float2 dy; // derivative of value with respect to screen-space y
};
extern "C" __device__ void tex_lookup_deriv_float3_2d(
float result[3],
Texture_handler_base const *self_base,
unsigned texture_idx,
tct_deriv_float2 const *coord,
Tex_wrap_mode const wrap_u,
Tex_wrap_mode const wrap_v,
float const crop_u[2],
float const crop_v[2]);
extern "C" __device__ void tex_lookup_deriv_float4_2d(
float result[4],
Texture_handler_base const *self_base,
unsigned texture_idx,
tct_deriv_float2 const *coord,
Tex_wrap_mode const wrap_u,
Tex_wrap_mode const wrap_v,
float const crop_u[2],
float const crop_v[2]);

The example texture runtime uses the CUDA tex2DGrad() functions to perform anisotropic hardware filtering. The derivatives provided to the texture lookup handlers can directly be used for those functions. The mipmaps are generated by Material_gpu_context::prepare_texture() in example_cuda_shared.h via the IImage_api::create_mipmaps() function.

[Previous] [Up] [Next]