C back end: incorrect code for calls to same function in different basic blocks

[richardmembarth]

We generate incorrect code for the following example:

extern "thorin" {
    fn cuda(int, (int, int, int), (int, int, int), fn() -> ()) -> ();
    fn bitcast[D, S](S) -> D;
}

fn nan() -> f32 { bitcast[f32](0x7fffffffi32) }

fn main(out: &[f32], idx: i32) -> () {
    cuda(0, (1, 1, 1), (1, 1, 1), || {
        if idx < 1 {
            out(idx) = nan();
            return()
        }
        out(idx) = nan();
    });
}

Generated CUDA code:

__global__ void lambda_crit_520(float* _523_559, int _524_560) {
    bool _562;
    _562 = _524_560 < 1;
    float* idx_564;
    idx_564 = _523_559 + _524_560;
    if (_562) goto l563; else goto l568;
    l563: ;
        float _566;
        union { float dst; int src; } u_566;
        u_566.src = 2147483647;
        _566 = u_566.dst;
        *idx_564 = _566;
        return ;
    l568: ;
        *idx_564 = _566;
        return ;
}

Problem: _566 is not set when entering l568.

[richardmembarth]

The actual problem are the two identical bitcasts which are mapped to the same thorin node and represent a constant primop:

lambda_crit_520(mem mem_521, fn(mem) return_522, [pf32]* _523, qs32 _524)
    bool _525 = lt _524, qs32 1
    pf32* idx_527 = lea _523, _524
    br_498(_525, if_then_526, if_else_531)

    if_else_531()
        mem _532 = store mem_521, idx_527, (pf32 bitcast qs32 2147483647)
        return_522(_532)

    if_then_526()
        mem _530 = store mem_521, idx_527, (pf32 bitcast qs32 2147483647)
        return_522(_530)

br_498(bool br_499, fn() br_500, fn() br_501)

For the first instance we generate correct code, but use the cached variable for the second. Our llvm-based back ends fold the bitcast (done by llvm) and we get directly a constant in both cases.