Disable vectorization for float8

src/targets/gpu/compile_gen.cpp

@@ -54,6 +54,11 @@ vectorize vectorize::elements(std::size_t axis,
                               const std::vector<shape>& inputs,
                               const std::vector<std::size_t>& sizes)
 {
+    // disable vectorization for fp8 types
+    if(std::any_of(inputs.begin(), inputs.end(), [&](auto ishape) {
+           return ishape.type() == migraphx::shape::float8_type;
+       }))
+        return {1, axis};
     if(std::all_of(
            inputs.begin(), inputs.end(), [&](const auto& s) { return s.lens()[axis] == 1; }))
         return {1, axis};
@@ -86,6 +91,11 @@ vectorize vectorize::elements(std::size_t axis,
 
 vectorize vectorize::elements(context& ctx, std::size_t axis, const std::vector<shape>& inputs)
 {
+    // disable vectorization for fp8 types
+    if(std::any_of(inputs.begin(), inputs.end(), [&](auto ishape) {
+           return ishape.type() == migraphx::shape::float8_type;
+       }))
+        return {1, axis};
     if(inputs.empty())
         return {1, axis};
     std::size_t n = std::max_element(inputs.begin(),
@@ -305,7 +315,7 @@ std::string generate_reduce(const module& m, const std::string& name)
             std::transform(
                 params.begin(), params.end(), params.begin(), [](auto s) { return "auto " + s; });
             return interpolate_string(inner_template,
-                                      {{"inner", inner_name},
+                                             {{"inner", inner_name},
                                        {"params", join_strings(params, ", ")},
                                        {"args", join_strings(args, ", ")},
                                        {"call", call_function}});

src/targets/gpu/kernels/include/migraphx/kernels/vectorize.hpp

@@ -237,9 +237,7 @@ template <index_int N, index_int Axis, class T>
 __device__ __host__ auto vectorize_tensor(T x)
 {
     constexpr auto shape = get_shape_c<T>{};
-    if constexpr(is_same<typename T::type, migraphx_fp8::fp8e4m3fnuz>{})
-        return x;
-    else if constexpr(shape.lens[Axis] == 1)
+    if constexpr(shape.lens[Axis] == 1)
         return x;
     else if constexpr(shape.strides[Axis] == 0)
         return tensor_step<N>(x, _c<Axis>);