Add missing tests and fixes from 4647

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

@@ -360,6 +360,12 @@ constexpr auto make_const_array(T x, Ts... xs)
     return integral_const_array<typename T::value_type, x, xs...>{};
 }
 
+template <class T, class N, class F>
+constexpr auto generate_const_array(N n, F f)
+{
+    return sequence_c<n>([=](auto... is) { return make_const_array(f(is)...); });
+}
+
 template <class T, class N, class F>
 constexpr auto generate_array(N n, F f)
 {

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

@@ -1,7 +1,7 @@
 /*
  * The MIT License (MIT)
  *
- * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved.
+ * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved.
  *
  * Permission is hereby granted, free of charge, to any person obtaining a copy
  * of this software and associated documentation files (the "Software"), to deal
@@ -206,6 +206,14 @@ MIGRAPHX_HIP_NORETURN inline __host__ __device__ void assert_fail(const source_l
 #define MIGRAPHX_CHECK(cond) \
     MIGRAPHX_ASSERT_FAIL(cond, #cond, __FILE__, __LINE__, __PRETTY_FUNCTION__)
 
+#ifdef CPPCHECK
+// NOLINTNEXTLINE
+#define MIGRAPHX_CAPTURE_SOURCE_LOCATION(T) T
+#define MIGRAPHX_ASSUME assert(cond)
+#define MIGRAPHX_UNREACHABLE assert(false)
+#define MIGRAPHX_ASSERT(cond) assert(cond)
+#define MIGRAPHX_WARN(cond, ...) assert(cond)
+#else
 #ifdef MIGRAPHX_DEBUG
 // NOLINTNEXTLINE
 #define MIGRAPHX_CAPTURE_SOURCE_LOCATION(T) source_location_capture<T>
@@ -221,6 +229,7 @@ MIGRAPHX_HIP_NORETURN inline __host__ __device__ void assert_fail(const source_l
 #define MIGRAPHX_ASSERT(cond)
 #define MIGRAPHX_WARN(...)
 #endif
+#endif
 
 #define MIGRAPHX_STATIC_ASSERT_FOR(...) \
     static_assert(__VA_ARGS__);         \

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

@@ -2,7 +2,7 @@
  *
  * The MIT License (MIT)
  *
- * Copyright (C) 2015-2025 Advanced Micro Devices, Inc. All rights reserved.
+ * Copyright (C) 2015-2026 Advanced Micro Devices, Inc. All rights reserved.
  *
  * Permission is hereby granted, free of charge, to any person obtaining a copy
  * of this software and associated documentation files (the "Software"), to deal
@@ -59,7 +59,7 @@ class numeric_limits;
 template <migraphx::fp8::f8_type T = migraphx::fp8::f8_type::fp8, bool FNUZ = true>
 struct float8
 {
-    uint8_t data;
+    uint8_t data = 0;
     // default constructor
     __device__ constexpr float8() = default;
     // default copy constructor
@@ -140,7 +140,7 @@ struct float8
            migraphx::fp8::rounding_mode rm = migraphx::fp8::rounding_mode::standard,
            uint32_t rng                    = 0)
     {
-        if(__builtin_is_constant_evaluated() or !FNUZ)
+        if(__builtin_is_constant_evaluated() or not FNUZ)
         {
             if constexpr(T == migraphx::fp8::f8_type::fp8)
             {
@@ -249,7 +249,7 @@ struct float8
     // upcast using device specific intrinsic
     constexpr __device__ operator float() const
     {
-        if(__builtin_is_constant_evaluated() or !FNUZ)
+        if(__builtin_is_constant_evaluated() or not FNUZ)
         {
             if constexpr(T == migraphx::fp8::f8_type::fp8)
             {
@@ -261,7 +261,7 @@ struct float8
         else
         {
             float fval      = 0;
-            uint32_t i32val = static_cast<uint32_t>(data);
+            uint32_t i32val = data;
 
             // upcast
             if constexpr(T == migraphx::fp8::f8_type::fp8)
@@ -312,7 +312,7 @@ struct float8
         }
         else
         {
-            if(T == migraphx::fp8::f8_type::bf8)
+            if constexpr(T == migraphx::fp8::f8_type::bf8)
             {
                 return (data == 0x7D) or (data == 0x7E) or (data == 0x7F) or (data == 0xFD) or
                        (data == 0xFE) or (data == 0xFF);
@@ -333,7 +333,7 @@ struct float8
         }
         else
         {
-            if(T == migraphx::fp8::f8_type::bf8)
+            if constexpr(T == migraphx::fp8::f8_type::bf8)
             {
                 return (data == 0x7C) or (data == 0xFC);
             }
@@ -370,16 +370,12 @@ struct float8
 
     __device__ constexpr bool operator<(const float8& rhs) const
     {
-        const auto we   = static_cast<float>(*this);
-        const auto them = static_cast<float>(rhs);
-        return we < them;
+        return static_cast<float>(*this) < static_cast<float>(rhs);
     }
 
     __device__ constexpr bool operator>(const float8& rhs) const
     {
-        const auto we   = static_cast<float>(*this);
-        const auto them = static_cast<float>(rhs);
-        return we > them;
+        return static_cast<float>(*this) > static_cast<float>(rhs);
     }
 };
 

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

@@ -118,6 +118,7 @@ __device__ constexpr uint8_t cast_to_f8(T f_x, bool stoch = false, uint32_t rng
     if(x == 0)
         return 0;
     // handle negative zero
+    // cppcheck-suppress compareValueOutOfTypeRangeError
     else if((sizeof(T) == 4 and x == 0x80000000) or (sizeof(T) == 2 and x == 0x8000))
     {
         return NegativeZeroNan ? 0 : 0x80; // For FNUZ types neg zero is just positive zero

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

@@ -1,7 +1,7 @@
 /*
  * The MIT License (MIT)
  *
- * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved.
+ * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved.
  *
  * Permission is hereby granted, free of charge, to any person obtaining a copy
  * of this software and associated documentation files (the "Software"), to deal
@@ -131,6 +131,9 @@ struct is_integral_constant<integral_constant<T, V>> : true_type
 template <index_int N>
 using index_constant = integral_constant<index_int, N>;
 
+template <index_int N>
+static constexpr auto index_c = index_constant<N>{};
+
 template <auto V>
 static constexpr auto _c = integral_constant<decltype(V), V>{}; // NOLINT
 

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

@@ -98,13 +98,12 @@ template <index_int N>
 constexpr auto slice_group()
 {
     return slice_size_transform{[](auto input, auto s) {
-        auto r = return_array_c([] {
+        return return_array_c([] {
             auto lens = decltype(s){}.lens.base();
             lens.back() *= N;
             lens -= 1;
             return decltype(input){}.lens.carry(lens) + index_int{1};
         });
-        return r;
     }};
 }
 

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

@@ -39,7 +39,7 @@ constexpr bool has_nonzero(index_ints<Ps...>)
     return ((Ps != 0) or ...);
 }
 
-template <index_int NTiles, class TileLens, class OutputShape, class Padding = index_ints<>>
+template <index_int NTiles, class TileLens, class OutputShape, class Padding = index_ints<0>>
 struct spatial_tiler
 {
     static constexpr auto keep_spatial()
@@ -72,39 +72,50 @@ struct spatial_tiler
     static constexpr index_int tiles_total() { return tiles_per_dim().product(); }
     static constexpr auto ndim() { return out_spatial_lens().size(); }
 
-    static constexpr bool is_padded()
+    static constexpr auto get_padding()
     {
-        return (out_spatial_lens() != tiles_per_dim() * output_lens());
+        if constexpr(Padding{}.size() < 2)
+        {
+            auto pre = transform(TileLens{}, [](auto) { return index_c<0>; });
+            return join(pre, pre);
+        }
+        else
+        {
+            return Padding{};
+        }
     }
 
-    static constexpr bool has_conv_padding() { return has_nonzero(Padding{}); }
-
     // Left (begin) padding per dim: (0, 0, left_h, left_w)
     static constexpr auto left_padding()
     {
-        return return_array_c([] {
-            constexpr auto p  = Padding{};
-            constexpr auto ns = p.size() / 2;
-            auto result       = array<index_int, ns + 2>(index_int{0});
-            for(index_int i = 0; i < ns; i++)
-                result[i + 2] = p[i];
-            return result;
+        constexpr auto p  = get_padding();
+        constexpr auto ns = p.size() / 2;
+        return generate_const_array<index_int>(_c<ns + 2>, [&](auto i) {
+            if constexpr(i < 2)
+                return index_c<0>;
+            else
+                return index_c<p[i - 2]>;
         });
     }
 
     // Total (left+right) padding per dim: (0, 0, left_h+right_h, left_w+right_w)
     static constexpr auto total_padding()
     {
-        return return_array_c([] {
-            constexpr auto p  = Padding{};
-            constexpr auto ns = p.size() / 2;
-            auto result       = array<index_int, ns + 2>(index_int{0});
-            for(index_int i = 0; i < ns; i++)
-                result[i + 2] = p[i] + p[i + ns];
-            return result;
+        constexpr auto p  = get_padding();
+        constexpr auto ns = p.size() / 2;
+        return generate_const_array<index_int>(_c<ns + 2>, [&](auto i) {
+            if constexpr(i < 2)
+                return index_c<0>;
+            else
+                return index_c<p[i - 2] + p[i - 2 + ns]>;
         });
     }
 
+    static constexpr bool is_padded()
+    {
+        return (out_spatial_lens() != (tiles_per_dim() * output_lens() + total_padding()));
+    }
+
     index idx;
     array<index_int, ndim()> tile_origin;
 
@@ -114,19 +125,10 @@ struct spatial_tiler
     static constexpr auto halo_lens_for()
     {
         constexpr auto halo_extra = [] {
-            if constexpr(has_conv_padding())
-            {
-                return return_array_c([] {
-                    return make_slice(InputShape{}, keep_spatial()).lens - out_spatial_lens() +
-                           total_padding();
-                });
-            }
-            else
-            {
-                constexpr auto input_spatial = make_slice(InputShape{}, keep_spatial()).lens;
-                return transform(
-                    input_spatial, out_spatial_lens(), [](auto is, auto os) { return is - os; });
-            }
+            return return_array_c([] {
+                return make_slice(InputShape{}, keep_spatial()).lens - out_spatial_lens() +
+                       total_padding();
+            });
         }();
         return transform(output_lens(), halo_extra, [](auto o, auto h) { return o + h; });
     }
@@ -167,19 +169,14 @@ struct spatial_tiler
         idx.local_stride(_c<hl.product()>, [&](auto i) {
             auto halo_multi = halo_shape.multi(i);
             auto src_pos    = tile_origin + halo_multi;
-            if constexpr(has_conv_padding())
+            auto input_pos  = src_pos - left_padding();
+            if constexpr(is_padded())
             {
-                constexpr auto pad = left_padding();
-                auto input_pos     = src_pos - pad;
                 smem[i] = in_bounds(input_pos, input_spatial) ? type{input_ch[input_pos]} : type{0};
             }
-            else if constexpr(is_padded())
-            {
-                smem[i] = in_bounds(src_pos, input_spatial) ? type{input_ch[src_pos]} : type{0};
-            }
             else
             {
-                smem[i] = input_ch[src_pos];
+                smem[i] = input_ch[input_pos];
             }
         });
 
@@ -203,7 +200,7 @@ struct spatial_tiler
     }
 };
 
-template <index_int NTiles, class TileLens, class OutputShape, class Padding = index_ints<>>
+template <index_int NTiles, class TileLens, class OutputShape, class Padding = index_ints<0>>
 __device__ auto make_spatial_tiler(index idx, TileLens, OutputShape, Padding = {})
 {
     using tiler_type = spatial_tiler<NTiles, TileLens, OutputShape, Padding>;

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

@@ -1,7 +1,7 @@
 /*
  * The MIT License (MIT)
  *
- * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved.
+ * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved.
  *
  * Permission is hereby granted, free of charge, to any person obtaining a copy
  * of this software and associated documentation files (the "Software"), to deal
@@ -28,6 +28,7 @@
 #include <migraphx/kernels/index.hpp>
 #include <migraphx/kernels/functional.hpp>
 #include <migraphx/kernels/tensor_view.hpp>
+#include <migraphx/kernels/uninitialized_buffer.hpp>
 #include <migraphx/kernels/copy.hpp>
 
 namespace migraphx {
@@ -61,8 +62,8 @@ struct tile
                 using type          = typename T::type;
                 constexpr auto s    = pad_shape(make_packed_shape(get_shape_c<T>{}));
                 constexpr auto size = s.element_space();
-                __shared__ type buffer[size];
-                auto b = make_tensor_view(buffer, s);
+                __shared__ uninitialized_buffer<type, size> buffer;
+                auto b = make_tensor_view(buffer.data(), s);
                 local_tensor_copy(idx, x, b);
                 f(b);
             };
@@ -77,8 +78,8 @@ struct tile
                 using type          = typename T::type;
                 constexpr auto s    = pad_shape(make_packed_shape(get_shape_c<T>{}));
                 constexpr auto size = s.element_space();
-                __shared__ type buffer[size];
-                auto b = make_tensor_view(buffer, s);
+                __shared__ uninitialized_buffer<type, size> buffer;
+                auto b = make_tensor_view(buffer.data(), s);
                 f(b);
                 local_tensor_copy(idx, b, x);
             };