Cppcheck fixes (#1195)

Fixes the #error when using cppcheck. This no longer suppresses cppcheck errors when including those errors. This fixes the cppcheck errors that was there already.

CMakeLists.txt

@@ -216,6 +216,7 @@ rocm_enable_cppcheck(
         CPPCHECK=1
         __device__=
         __host__=
+        __global__=
 )
 
 enable_testing()

src/api/include/migraphx/migraphx.hpp

@@ -314,6 +314,7 @@ struct interface_base : Base
                 T** y = reinterpret_cast<T**>(out);
                 T* x  = reinterpret_cast<T*>(input);
                 assert(x != nullptr and y != nullptr and *y == nullptr);
+                // cppcheck-suppress useSmartPointer
                 *y = new T(*x); // NOLINT
             });
         };
@@ -339,6 +340,7 @@ struct interface_base : Base
     template <class T, class Setter, class F>
     void set_auto_fp(Setter setter, F f)
     {
+        // cppcheck-suppress constParameter
         return set_fp<T>(setter, [=](T& obj, auto out, auto... xs) {
             auto_invoke(f, out, obj, auto_convert_param(rank<2>{}, xs)...);
         });

src/driver/marker_roctx.cpp

@@ -17,7 +17,7 @@ class marker_roctx
     std::function<int(const char*)> sym_roctx_range_push;
     std::function<int()> sym_roctx_range_pop;
 
-    uint64_t range_id;
+    uint64_t range_id = 0;
 
     public:
     marker_roctx()

src/include/migraphx/filesystem.hpp

@@ -3,7 +3,10 @@
 
 #include <migraphx/config.hpp>
 
-#if defined(__has_include) && !defined(CPPCHECK)
+#if defined(CPPCHECK)
+#define MIGRAPHX_HAS_FILESYSTEM 1
+#define MIGRAPHX_HAS_FILESYSTEM_TS 1
+#elif defined(__has_include)
 #if __has_include(<filesystem>) && __cplusplus >= 201703L
 #define MIGRAPHX_HAS_FILESYSTEM 1
 #else

src/include/migraphx/optional.hpp

@@ -3,7 +3,10 @@
 
 #include <migraphx/config.hpp>
 
-#if defined(__has_include) && !defined(CPPCHECK)
+#if defined(CPPCHECK)
+#define MIGRAPHX_HAS_OPTIONAL 1
+#define MIGRAPHX_HAS_OPTIONAL_TS 1
+#elif defined(__has_include)
 #if __has_include(<optional>) && __cplusplus >= 201703L
 #define MIGRAPHX_HAS_OPTIONAL 1
 #else

src/onnx/parse_mean.cpp

@@ -24,13 +24,16 @@ struct parse_mean : op_parser<parse_mean>
         auto divisor = info.add_literal(
             migraphx::literal{migraphx::shape{args[0]->get_shape().type()}, {num_data}});
 
-        return std::accumulate(args.begin(), args.end(), args[0], [&](auto& mean, auto& data_i) {
-            // Pre-divide each tensor element-wise by n to reduce risk of overflow during summation
-            data_i = info.add_broadcastable_binary_op("div", data_i, divisor);
-
-            if(data_i != args[0])
-                return info.add_broadcastable_binary_op("add", mean, data_i);
-            return data_i;
+        // TODO: Only divide when using floating-point
+        return std::accumulate(args.begin() + 1,
+                               args.end(),
+                               info.add_broadcastable_binary_op("div", args[0], divisor),
+                               [&](auto mean, auto data_i) {
+                                   // Pre-divide each tensor element-wise by n to reduce risk of
+                                   // overflow during summation
+                                   auto div =
+                                       info.add_broadcastable_binary_op("div", data_i, divisor);
+                                   return info.add_broadcastable_binary_op("add", mean, div);
                                });
     }
 };

src/simplify_algebra.cpp

@@ -995,7 +995,7 @@ struct find_split_transpose
         auto axis = any_cast<op::slice>(slc->get_operator()).axes.front();
         auto it   = std::find(perm.begin(), perm.end(), axis);
         assert(it != perm.end());
-        auto axis_new = static_cast<int64_t>(std::distance(perm.begin(), it));
+        int64_t axis_new = std::distance(perm.begin(), it);
 
         for(auto in : split_outputs)
         {

src/targets/gpu/compile_hip.cpp

@@ -133,6 +133,7 @@ struct hiprtc_program
         std::vector<char> buffer(n);
         MIGRAPHX_HIPRTC(hiprtcGetProgramLog(prog.get(), buffer.data()));
         assert(buffer.back() == 0);
+        // cppcheck-suppress returnDanglingLifetime
         return {buffer.begin(), buffer.end() - 1};
     }
 

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

@@ -118,15 +118,13 @@ constexpr roalign_settings<Ts...> make_roalign_settings(Ts... xs)
 }
 
 template <class T, class U, class V, class W, class Settings>
-__device__ void roialign(const T& x_t, const U& rois_t, const V& ind_t, const W& y_t, Settings s)
+__device__ void roialign(const T& x_t, const U& rois_t, const V& ind_t, W& y_t, Settings s)
 {
     auto index      = make_index();
     const auto x    = x_t.begin();
     const auto rois = rois_t.begin();
     const auto ind  = ind_t.begin();
 
-    auto out_ptr = y_t.begin();
-
     // input shape
     auto x_lens      = x_t.get_shape().lens;
     auto channel_num = x_lens[1];
@@ -176,7 +174,7 @@ __device__ void roialign(const T& x_t, const U& rois_t, const V& ind_t, const W&
         const auto offset_x = x + ((batch_ind * channel_num + c) * in_dims[0] * in_dims[1]);
         if constexpr(s.is_avg_pooling)
         {
-            out_ptr[i] = calc_pooling(offset_x,
+            y_t[i] = calc_pooling(offset_x,
                                   roi_starts,
                                   bin_size,
                                   {ph, pw},
@@ -187,7 +185,7 @@ __device__ void roialign(const T& x_t, const U& rois_t, const V& ind_t, const W&
         }
         else
         {
-            out_ptr[i] = calc_pooling(offset_x,
+            y_t[i] = calc_pooling(offset_x,
                                   roi_starts,
                                   bin_size,
                                   {ph, pw},