Dropping arguments in inlining #648
Comments
|
Thanks for the report! That's a concerning bug. Arguments must not be dropped in entrypoint kernels. I think this should be an easy fix. Here's a smaller reproducer: knl = lp.make_kernel(
"{ : }",
"""
a[0] = 1
""",
[lp.GlobalArg("a,b,c,d,e",
shape=(10,),
dtype="float64")])
print(lp.generate_code_v2(knl).device_code())which generates the kernel: __kernel void __attribute__ ((reqd_work_group_size(1, 1, 1))) loopy_kernel(__global double *__restrict__ a)
{
a[0] = 1.0;
}Notice how the arguments |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Hi. I noticed an issue with the argument passing that has been introduced in #631 (I believe).
When we inline an inner kernel in a wrapper kernel, where the inner kernels takes many arguments but the code only depends on a subset of it, then in the C code, the inlined kernel only takes a subset of the arguments as parameters. A minimal example can be found here https://gist.github.com/sv2518/d6690f63b6827750fdaed9b01b4d9cea. I noticed that this only happens to
CTargets, not for whatever Loo.py's default target is.This is problematic in Firedrake for the vectorisation where we generate a loopy kernel first, do the inlining and vectorisation transformations and C codegen after that, and then pass the arguments dynamically. Meaning the arguments for the kernel before the transformations must be the same as after it. The MFE I linked above is coming from our test suite, the corresponding test is
tests/regression/test_par_loops.py::test_dict_order_parallel.Also sorry for the noise recently, it's mostly because I am still trying to get the Firedrake vectorisation PR landed.
The text was updated successfully, but these errors were encountered: