GPU Codegen bug: bad cast of input scalar

[FlorianDeconinck]

Dear DaCe team,

We encountered a bug in the codegen (SDFG linked and repro' below).

The quirk is that a tasklet takes a np.float32 scalar but is given a np.float64 - this triggers a malformed cast inside the kernel on the callsite of a nested function, e.g.

void nested_sdfg_0_0_3(const double&  scalar, float* __restrict__ out_field, int __i, int __j);

__global__ void  __launch_bounds__(32) horizontal_loop_139965476328208_0_0_5(float * __restrict__ stencil_out, const double stencil_scalar) {
       #code removed
       nested_sdfg_0_0_3(*(float *)(&stencil_scalar), &stencil_out[0], __i, __j);
       #--------------------------------^ bad idea
}

Rep:

import dace

sdfg = dace.SDFG.from_file("./bad_cast_on_scalar.sdfgz.zip")
sdfg.compile()

[romanc]

according to @ThrudPrimrose

This can be detected at the SDFG level, and in my opinion it should be handled at the SDFG level.

A pass can be written that detects if outside data descriptor is connecting to an nsdfg that expects a different dtype for the connector. A typecast tasklet can be inserted to change what is passed to the NSDFG. (And the pass can adapt to whether an array is being passed or a scalar).

Handling it safely just in codegen is going to be hard. If the data being passed is not constant within the nsdfg, then we need to cast it back and assign to the original value. Which seems like much better done in a transformation.

If having C++ code in a tasklet in that stage is not desirable, we can define sympy user functions that handle the cast that are correctly lowered, such that all tasklets remain Python. For this a small discussion might be good with other DaCe core team, e.g. similar to int_floor we can have dace.float32(val) as a function that converts input to fp32. This was done by the Fortran frontend.

I think for the general improvement of QoL, I would propose we integrate it as a pre-code generation pass and make "C-style" casts into an error. (Again needs discussion). In my experience avoiding Cpp tasklets is just better, so we might define float32 etc. as sympy user functions that we lower to static_cast later at codegen.