Broadcast after concat runtime error

Issue

I encountered an error when applying the concat operation and then the broadcast operation. This occurred when making a [2] shape parameter, concatenating a NaN to make a [3] shape, and then trying to broadcast to a [3, 2] shape.

[ RUN ] test_isnan_broadcast
In file included from main.cpp:3:
./migraphx/kernels/pointwise.hpp:21:9: error: static_assert failed due to requirement 'vec_size<bool>() == 2U' "Vector mismatch size"
static_assert(vec_size<T>() == N, "Vector mismatch size");
^ ~~~~~~~~~~~~~~~~~~
./migraphx/kernels/pointwise.hpp:44:30: note: in instantiation of function template specialization 'migraphx::implicit_conversion_op<bool>::operator unsigned char __attribute__((ext_vector_type(2)))<2, unsigned char>' requested here
out[multi_idx] = implicit_conversion(f(ps[multi_idx]...));
^
./migraphx/kernels/index.hpp:39:13: note: in instantiation of function template specialization 'migraphx::pointwise_tensor(migraphx::index, (lambda at main.cpp:22:15), migraphx::tensor_view<unsigned char __attribute__((ext_vector_type(2))), migraphx::shape<migraphx::integral_const_array<unsigned int, 3, 1>, migraphx::integral_const_array<unsigned int, 1, 1>>>, migraphx::tensor_view<float, migraphx::shape<migraphx::integral_const_array<unsigned int, 3, 1>, migraphx::integral_const_array<unsigned int, 1, 0>>>)::(anonymous class)::operator()(migraphx::tensor_view<float, migraphx::shape<migraphx::integral_const_array<unsigned int, 3, 1>, migraphx::integral_const_array<unsigned int, 1, 0>>>)::(anonymous class)::operator()<unsigned int>' requested here
f(i);
^
./migraphx/kernels/pointwise.hpp:42:13: note: in instantiation of function template specialization 'migraphx::index::global_stride<(lambda at ./migraphx/kernels/pointwise.hpp:42:55)>' requested here
idx.global_stride(out.get_shape().elements(), [&](auto i) {
^
./migraphx/kernels/pointwise.hpp:55:9: note: in instantiation of function template specialization 'migraphx::pointwise_tensor<(lambda at main.cpp:22:15), migraphx::tensor_view<unsigned char __attribute__((ext_vector_type(2))), migraphx::shape<migraphx::integral_const_array<unsigned int, 3, 1>, migraphx::integral_const_array<unsigned int, 1, 1>>>, migraphx::tensor_view<float, migraphx::shape<migraphx::integral_const_array<unsigned int, 3, 1>, migraphx::integral_const_array<unsigned int, 1, 0>>>>' requested here
pointwise_tensor(idx, f, xs...);
^
main.cpp:22:5: note: in instantiation of function template specialization 'migraphx::pointwise<(lambda at main.cpp:22:15), void, void>' requested here
pointwise(MIGRAPHX_LIFT(main_pointwise0), private_p0,private_p1);
^
In file included from main.cpp:3:
./migraphx/kernels/pointwise.hpp:22:16: error: first argument to __builtin_convertvector must be a vector
return __builtin_convertvector(x, vec<U, N>);
^ ~
2 errors generated when compiling for gfx1030.
module: "main"
output = @param:output -> bool_type, {3, 2}, {2, 1}
main:@0 = @literal{nan} -> float_type, {1}, {0}
x = @param:x -> float_type, {2}, {1}
main:@1 = hip::allocate[shape=float_type, {3}, {1},tag=] -> float_type, {3}, {1}
main:@2 = gpu::concat[axis=0](x,main:@0,main:@1) -> float_type, {3}, {1}
main:@3 = broadcast[axis=0,out_lens={3, 2}](main:@2) -> float_type, {3, 2}, {1, 0}
main:@4 = gpu::precompile_op[op=pointwise](main:@3,output), [main:pointwise0] -> bool_type, {3, 2}, {2, 1}
module: "main:pointwise0"
x0 = @param:x0 -> float_type, {1}, {0}
main:pointwise0:@0 = isnan(x0) -> bool_type, {1}, {0}
main:pointwise0:@1 = @return(main:pointwise0:@0)

FAILED: test_isnan_broadcast
what(): /code/AMDMIGraphX/src/compile_src.cpp:41: compile: Output file missing: main.o

You can see the test_isnan_broadcast here: https://github.com/ROCmSoftwarePlatform/AMDMIGraphX/commit/bf9456e01ed5f8b135028c67d71522ade8601a01.
Reversing the order of the operations (broadcast -> concat) produces no issue.

Notes from Paul

There might be a bug in the pointwise kernel, but we should be optimizing this anyways in simplify_algebra(and we dont)
find_inner_broadcast only works on binary operators, we should extend it for unary operators as well

A quick fix is just to skip vectorization when all the inputs are broadcasted
You can update auto_vectorize_impl like this:

template <class F, class... Ts>
inline __device__ __host__ auto auto_vectorize_impl(F f, Ts... xs)
{
    // TODO: Just check there a single axis of 1
    constexpr bool packed_or_broadcasted =
        ((xs.get_shape().packed() or xs.get_shape().broadcasted()) and ...);
    constexpr bool all_broadcasted =
        (xs.get_shape().broadcasted() and ...);
    if constexpr(packed_or_broadcasted and not all_broadcasted)
    {
        ...
    }
    else
    {
        f(xs...);
    }
}

Start by reproducing test_isnan_broadcast from commit bf9456e01ed5f8b135028c67d71522ade8601a01 and read migraphx/kernels/pointwise.hpp, especially auto_vectorize_impl. Compare the concat-then-broadcast and broadcast-then-concat shapes, then use the issue's suggested behavior as the completion target and confirm the test compiles successfully.