Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Constant emitter causes kernel recompilation with geometry change #1033

Open
futscdav opened this issue Jan 10, 2024 · 2 comments
Open

Constant emitter causes kernel recompilation with geometry change #1033

futscdav opened this issue Jan 10, 2024 · 2 comments

Comments

@futscdav
Copy link

Summary

When updating the geometry of the scene and using constant emitter, the jit kernels will be recompiled. Swapping to e.g. spot emitter will not recompile the kernel, as expected. Here's a sample repro and diff of the kernels. I'm not 100% certain if this is caused strictly by changing the geometry, but changing e.g. bsdf reflectance data does not result in different kernel.

import drjit as dr
import mitsuba as mi
import numpy as np

mi.set_variant("cuda_ad_rgb")

dr.set_log_level(dr.LogLevel.Info)
dr.set_flag(dr.JitFlag.PrintIR, False)

vert_count = 3
face_count = 1

verts = np.array([
    [-1, 0, 0], [1, 0, 0], [1, -1, 0]
]).astype(np.float32)
faces = np.array([
    2, 1, 0
])

v_array = mi.Float(np.ravel(verts))

mesh = mi.Mesh(
    'mesh',
    vertex_count=vert_count,
    face_count=face_count,
    has_vertex_normals=False,
    has_vertex_texcoords=False,
)
mesh_params = mi.traverse(mesh)
mesh_params['vertex_positions'] = v_array
mesh_params['faces'] = np.ravel(faces)
mesh_params.update()

scene = mi.load_dict({
    'type': 'scene',
    'integrator': {'type': 'direct'},
    # 'emitter': {
    #     'type': 'spot',
    #     'intensity': {'type': 'rgb', 'value': 100.},
    #     'to_world': mi.ScalarTransform4f.look_at(
    #         [0, 0, 5], [0, 0, 0], [0, 1, 0]
    #     ),
    # },
    'emitter': {'type': 'constant'},
    'shape': mesh,
    'sensor': {
        'type': 'perspective',
        'to_world': mi.ScalarTransform4f.look_at(
            [0, 0, 5], [0, 0, 0], [0, 1, 0]
        ),
        'film': {
            'type': 'hdrfilm',
            'width': 1,
            'height': 1,
            'pixel_format': 'rgb',
            'sample_border': True,
        },
    },
})

params = mi.traverse(scene)
for i in range(3):
  print(f'Iteration {i}')
  v_array = v_array + 1e-6
  params['shape.vertex_positions'] = dr.detach(v_array)
  params.update()
  image = mi.render(scene, params, spp=4)

The kernel PTX is identical except for 3 values (addresses?) inside of VCall Emitter::sample_direction:

...
.visible .func (.param .align 4 .b8 result[36]) __direct_callable__ac5362b1627661b56ab644a4f6b0d49e(.reg .u32 self, .reg .u64 data, .param .align 4 .b8 params[20]) {
    // VCall: mitsuba::Emitter::sample_direction()
    .reg.b8   %b <43>; .reg.b16 %w<43>; .reg.b32 %r<43>;
    .reg.b64  %rd<43>; .reg.f32 %f<43>; .reg.f64 %d<43>;
    .reg.pred %p <43>;

    ld.global.f32 %f4, [data+0];
    mov.b32 %f5, 0x40000000;
    ld.param.f32 %f6, [params+16];
    neg.ftz.f32 %f7, %f6;
    mov.b32 %f8, 0x3f800000;
    fma.rn.ftz.f32 %f9, %f5, %f7, %f8;
    neg.ftz.f32 %f10, %f9;
    fma.rn.ftz.f32 %f11, %f9, %f10, %f8;
    mov.b32 %f12, 0x0;
    max.ftz.f32 %f13, %f11, %f12;
    sqrt.approx.ftz.f32 %f14, %f13;
    mov.b32 %f15, 0x40c90fdb;
    ld.param.f32 %f16, [params+12];
    mul.ftz.f32 %f17, %f15, %f16;
    cos.approx.ftz.f32 %f18, %f17;
    mul.ftz.f32 %f19, %f14, %f18;
    mov.b32 %f20, 0x3ebbb30b;
    ld.param.f32 %f21, [params+8];
    mov.b32 %f22, 0x3e41813c;  <-- different value 1
    sub.ftz.f32 %f23, %f21, %f22;
    ld.param.f32 %f24, [params+4];
    mov.b32 %f25, 0x3ebbc0f0;  <-- different value 2
    sub.ftz.f32 %f26, %f24, %f25;
    ld.param.f32 %f27, [params+0];
    mov.b32 %f28, 0x3f1fa4b9;  <-- different value 3
    sub.ftz.f32 %f29, %f27, %f28;
    mul.ftz.f32 %f30, %f29, %f29;
    fma.rn.ftz.f32 %f31, %f26, %f26, %f30;
...

System configuration

System information:

OS: Debian
CPU: x86_64
GPU: RTX 4090
Python version: 3.11
CUDA version: 12.0
NVidia driver: 525.147.05

@merlinND
Copy link
Member

Hi @futscdav,

Thanks for the report and investigation. Could these values correspond to floats related to the scene's bounding sphere radius and center? They seem to be computed as scalar quantities in constant.cpp.

@futscdav
Copy link
Author

That seems like a reasonable explanation, using fixed center and radius like

{
  'type': 'sphere',
  'to_world': mi.ScalarTransform4f.translate([0, 0, 0]).scale(-1000.0),
  'emitter': {'type': 'area', 'radiance': 1.0},
}

as a workaround works fine.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

No branches or pull requests

2 participants