Skip to content

Unrecoverable error on Bool comparisons #462

@seanremy

Description

@seanremy

Platform: Windows 11
drjit version: 1.2.0
backend: CUDA version 13.0

The following simple example:

import drjit as dr
from drjit.auto import Bool

dr.set_flag(dr.JitFlag.Debug, True)
Bool(True, True) > Bool(True, False)

yields the following error:

Dr.Jit encountered an unrecoverable error and will now shut
down. Please re-run your program in debug mode to check for
out-of-bounds reads, writes, and other sources of undefined
behavior. You can do so by calling

   dr.set_flag(dr.JitFlag.Debug, True)

at the beginning of the program. If these additional checks
fail to pinpoint the problem, then you have likely found a
bug. We are happy to help investigate and fix the problem if
you can you create a self-contained reproducer and submit it
at https://github.com/mitsuba-renderer/drjit.

The error message of this specific failure is as follows:
>>> jit_cuda_compile(): compilation failed. Please see the PTX assembly listing and error message below:

.version 9.0
.target sm_89
.address_size 64

.entry drjit_fc6c7ad2bf726887881cbcc668dc9424(.param .align 8 .b8 params[32]) {
    .pragma "enable_smem_spilling";

    .reg.b8  %b <7>; .reg.b16  %w<7>; .reg.b32 %r<7>;
    .reg.b64 %rd<7>; .reg.f16  %h<7>; .reg.f32 %f<7>;
    .reg.f64 %d <7>; .reg.pred %p<7>;

    mov.u32 %r0, %ctaid.x;
    mov.u32 %r1, %ntid.x;
    mov.u32 %r2, %tid.x;
    mad.lo.u32 %r0, %r0, %r1, %r2;
    ld.param.u32 %r2, [params];
    setp.ge.u32 %p0, %r0, %r2;
    @%p0 ret;

body: // sm_89
    ld.param.u64 %rd0, [params+8];
    mad.wide.u32 %rd0, %r0, 1, %rd0;
    ld.global.cs.b8 %w0, [%rd0];
    setp.ne.u16 %p4, %w0, 0;
    ld.param.u64 %rd0, [params+16];
    mad.wide.u32 %rd0, %r0, 1, %rd0;
    ld.global.cs.b8 %w0, [%rd0];
    setp.ne.u16 %p5, %w0, 0;
    setp.gt.pred %p6, %p4, %p5;
    ld.param.u64 %rd0, [params+24];
    mad.wide.u32 %rd0, %r0, 1, %rd0;
    selp.u16 %w0, 1, 0, %p6;
    st.global.cs.b8 [%rd0], %w0;
    ret;
}


ptxas application ptx input, line 29; error   : Unexpected instruction types specified for 'setp'
ptxas fatal   : Ptx assembly aborted due to errors

No error if I explicitly use the scalar backend. No error if I use single element arrays, like this:
Bool([True]) > Bool([False])

Metadata

Metadata

Assignees

No one assigned

    Labels

    No labels
    No labels

    Type

    No type

    Projects

    No projects

    Milestone

    No milestone

    Relationships

    None yet

    Development

    No branches or pull requests

    Issue actions