Merge branch 'develop' of github.com:ROCmSoftwarePlatform/AMDMIGraphX into dyn_conv

src/include/migraphx/op/step.hpp

@@ -24,16 +24,11 @@
 #ifndef MIGRAPHX_GUARD_OPERATORS_STEP_HPP
 #define MIGRAPHX_GUARD_OPERATORS_STEP_HPP
 
-#include "migraphx/stringutils.hpp"
-#include <array>
 #include <migraphx/check_shapes.hpp>
 #include <migraphx/argument.hpp>
-#include <migraphx/functional.hpp>
 #include <migraphx/config.hpp>
-#include <migraphx/lifetime.hpp>
+#include <migraphx/value.hpp>
 #include <migraphx/op/normalize_attribute.hpp>
-#include <cmath>
-#include <utility>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {

src/include/migraphx/op/sub.hpp

@@ -24,16 +24,9 @@
 #ifndef MIGRAPHX_GUARD_OPERATORS_SUB_HPP
 #define MIGRAPHX_GUARD_OPERATORS_SUB_HPP
 
-#include <array>
-#include <migraphx/op/binary.hpp>
-#include <migraphx/check_shapes.hpp>
-#include <migraphx/stringutils.hpp>
-#include <migraphx/streamutils.hpp>
-#include <migraphx/literal.hpp>
-#include <migraphx/shape_for_each.hpp>
 #include <migraphx/config.hpp>
+#include <migraphx/op/binary.hpp>
 #include <cmath>
-#include <utility>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {

src/include/migraphx/op/tan.hpp

@@ -24,16 +24,9 @@
 #ifndef MIGRAPHX_GUARD_OPERATORS_TAN_HPP
 #define MIGRAPHX_GUARD_OPERATORS_TAN_HPP
 
-#include <array>
-#include <migraphx/op/unary.hpp>
-#include <migraphx/check_shapes.hpp>
-#include <migraphx/stringutils.hpp>
-#include <migraphx/streamutils.hpp>
-#include <migraphx/literal.hpp>
-#include <migraphx/shape_for_each.hpp>
 #include <migraphx/config.hpp>
+#include <migraphx/op/unary.hpp>
 #include <cmath>
-#include <utility>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {

src/include/migraphx/op/tanh.hpp

@@ -24,16 +24,9 @@
 #ifndef MIGRAPHX_GUARD_OPERATORS_TANH_HPP
 #define MIGRAPHX_GUARD_OPERATORS_TANH_HPP
 
-#include <array>
-#include <migraphx/op/unary.hpp>
-#include <migraphx/check_shapes.hpp>
-#include <migraphx/stringutils.hpp>
-#include <migraphx/streamutils.hpp>
-#include <migraphx/literal.hpp>
-#include <migraphx/shape_for_each.hpp>
 #include <migraphx/config.hpp>
+#include <migraphx/op/unary.hpp>
 #include <cmath>
-#include <utility>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {

src/include/migraphx/op/transpose.hpp

@@ -24,14 +24,11 @@
 #ifndef MIGRAPHX_GUARD_OPERATORS_TRANSPOSE_HPP
 #define MIGRAPHX_GUARD_OPERATORS_TRANSPOSE_HPP
 
-#include <array>
 #include <migraphx/check_shapes.hpp>
 #include <migraphx/argument.hpp>
-#include <migraphx/functional.hpp>
 #include <migraphx/config.hpp>
-#include <migraphx/lifetime.hpp>
-#include <cmath>
-#include <utility>
+#include <migraphx/value.hpp>
+#include <migraphx/op/normalize_attribute.hpp>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {

src/include/migraphx/op/unary_not.hpp

@@ -24,10 +24,9 @@
 #ifndef MIGRAPHX_GUARD_OPERATORS_UNARY_NOT_HPP
 #define MIGRAPHX_GUARD_OPERATORS_UNARY_NOT_HPP
 
-#include <migraphx/op/unary.hpp>
-#include <migraphx/operation.hpp>
-#include <migraphx/check_shapes.hpp>
 #include <migraphx/config.hpp>
+#include <migraphx/op/unary.hpp>
+#include <cmath>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {

src/include/migraphx/op/unsqueeze.hpp

@@ -24,16 +24,11 @@
 #ifndef MIGRAPHX_GUARD_OPERATORS_UNSQUEEZE_HPP
 #define MIGRAPHX_GUARD_OPERATORS_UNSQUEEZE_HPP
 
-#include <array>
 #include <migraphx/check_shapes.hpp>
-#include <migraphx/stringutils.hpp>
-#include <migraphx/streamutils.hpp>
-#include <migraphx/shape_for_each.hpp>
+#include <migraphx/argument.hpp>
 #include <migraphx/config.hpp>
+#include <migraphx/value.hpp>
 #include <migraphx/op/normalize_attribute.hpp>
-#include <migraphx/lifetime.hpp>
-#include <cmath>
-#include <utility>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {

src/include/migraphx/op/where.hpp

@@ -24,18 +24,11 @@
 #ifndef MIGRAPHX_GUARD_OPERATORS_WHERE_HPP
 #define MIGRAPHX_GUARD_OPERATORS_WHERE_HPP
 
-#include <array>
-#include <migraphx/argument.hpp>
-#include <migraphx/par_for.hpp>
 #include <migraphx/check_shapes.hpp>
-#include <migraphx/stringutils.hpp>
-#include <migraphx/streamutils.hpp>
-#include <migraphx/shape_for_each.hpp>
+#include <migraphx/argument.hpp>
 #include <migraphx/config.hpp>
 #include <migraphx/value.hpp>
-#include <migraphx/op/normalize_attribute.hpp>
-#include <cmath>
-#include <utility>
+#include <migraphx/par_for.hpp>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {

src/include/migraphx/par_dfor.hpp

@@ -25,6 +25,7 @@
 #define MIGRAPHX_GUARD_RTGLIB_PAR_DFOR_HPP
 
 #include <migraphx/par_for.hpp>
+#include <migraphx/dfor.hpp>
 #include <migraphx/functional.hpp>
 #include <array>
 #include <numeric>

src/targets/gpu/compile_hip_code_object.cpp

@@ -51,9 +51,9 @@ static const char* const make_tensor_template = R"__migraphx__(
 template<>
 struct make_tensor<${n}>
 {
-    static __device__ auto apply(void* p)
+    static __device__ auto apply(void* __restrict__ p)
     {
-        return make_tensor_view(reinterpret_cast<${type}*>(p), make_shape(${lens}, ${strides}));
+        return make_tensor_view(reinterpret_cast<${type}* __restrict__>(p), make_shape(${lens}, ${strides}));
     }
 };
 )__migraphx__";

src/targets/gpu/include/migraphx/gpu/device_name.hpp

@@ -33,6 +33,8 @@ namespace gpu {
 
 std::string get_device_name();
 
+int get_device_id();
+
 } // namespace gpu
 } // namespace MIGRAPHX_INLINE_NS
 } // namespace migraphx

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

@@ -27,6 +27,7 @@
 #include <migraphx/kernels/hip.hpp>
 #include <migraphx/kernels/types.hpp>
 #include <migraphx/kernels/integral_constant.hpp>
+#include <migraphx/kernels/type_traits.hpp>
 
 namespace migraphx {
 
@@ -53,29 +54,51 @@ struct index
         return blockDim.x; // NOLINT
     }
 #endif
+    template <class N, class Stride>
+    static constexpr auto max_stride_iterations(N n, Stride stride)
+    {
+        return (n - _c<1>) / stride + _c<1>;
+    }
 
-    template <class F>
-    __device__ void global_stride(index_int n, F f) const
+    template <class F, class N, class Stride>
+    static constexpr void for_stride(index_int start, N n, Stride stride, F f)
     {
-        const auto stride = nglobal();
-        for(index_int i = global; i < n; i += stride)
+        if constexpr(not is_integral<N>{} and not is_integral<Stride>{} and
+                     max_stride_iterations(n, stride) == 1)
         {
-            f(i);
+            if constexpr(stride > n)
+            {
+                if(start < n)
+                    f(start);
+            }
+            else
+            {
+                f(start);
+            }
+        }
+        else
+        {
+            for(index_int i = start; i < n; i += stride)
+            {
+                f(i);
+            }
         }
     }
 
-    template <class F>
-    __device__ void local_stride(index_int n, F f) const
+    template <class F, class N>
+    __device__ void global_stride(N n, F f) const
     {
-        const auto stride = nlocal();
-        for(index_int i = local; i < n; i += stride)
-        {
-            f(i);
-        }
+        for_stride(global, n, nglobal(), f);
+    }
+
+    template <class F, class N>
+    __device__ void local_stride(N n, F f) const
+    {
+        for_stride(local, n, nlocal(), f);
     }
 };
 
-inline __device__ index make_index()
+inline __device__ __attribute__((const)) index make_index()
 {
     return index{blockIdx.x * blockDim.x + threadIdx.x, threadIdx.x, blockIdx.x}; // NOLINT
 }

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

@@ -186,7 +186,8 @@ __device__ auto auto_preload(index idx)
 {
     return make_transform([=](auto f, auto... xs) {
         auto invoke = [=](auto... ys) {
-            __syncthreads();
+            if constexpr((Bs or ...))
+                __syncthreads();
             f(ys...);
         };
         join(invoke, preload_copy<Bs>(idx, xs)...);

src/targets/gpu/target.cpp

@@ -53,6 +53,7 @@
 #include <migraphx/gpu/compile_ops.hpp>
 #include <migraphx/gpu/concat_gpu_opt.hpp>
 #include <migraphx/gpu/context.hpp>
+#include <migraphx/gpu/device_name.hpp>
 #include <migraphx/gpu/fuse_mlir.hpp>
 #include <migraphx/gpu/fuse_ops.hpp>
 #include <migraphx/gpu/prefuse_ops.hpp>
@@ -162,7 +163,7 @@ std::vector<pass> target::get_passes(migraphx::context& gctx, const compile_opti
 
 std::string target::name() const { return "gpu"; }
 
-migraphx::context target::get_context() const { return context{}; }
+migraphx::context target::get_context() const { return context(gpu::get_device_id()); }
 
 argument target::copy_to(const argument& arg) const { return gpu::to_gpu(arg); }
 

test/api/CMakeLists.txt

@@ -32,6 +32,17 @@ function(add_api_test TEST_NAME TEST_SRC TEST_DIR)
     add_dependencies(check ${NAME})
 endfunction()
 
+# Workaround: C file dont work with clang-tidy right now, need a fix in rocm-cmake
+function(add_c_api_test TEST_NAME TEST_SRC TEST_DIR)
+    set(NAME test_api_${TEST_NAME})
+    add_executable(${NAME} EXCLUDE_FROM_ALL ${TEST_SRC})
+    target_link_libraries(${NAME} migraphx_c migraphx)
+    target_include_directories(${NAME} PUBLIC ../include)
+    add_test(NAME ${NAME} COMMAND $<TARGET_FILE:${NAME}> WORKING_DIRECTORY ${TEST_DIR}) 
+    add_dependencies(tests ${NAME})
+    add_dependencies(check ${NAME})
+endfunction()
+
 add_api_test(array_base test_array_base.cpp ${TEST_ONNX_DIR})
 add_api_test(assign test_assign.cpp ${TEST_ONNX_DIR})
 add_api_test(compile_options test_compile_options.cpp ${TEST_ONNX_DIR})
@@ -40,6 +51,7 @@ add_api_test(module_construct test_module_construct.cpp ${TEST_ONNX_DIR})
 add_api_test(ref test_cpu.cpp ${TEST_ONNX_DIR})
 add_api_test(save_load test_save_load.cpp ${TEST_ONNX_DIR})
 add_api_test(op test_op_construct.cpp ${TEST_ONNX_DIR})
+add_c_api_test(c_op test_c_op_construct.c ${TEST_ONNX_DIR})
 add_api_test(custom_op test_custom_op.cpp ${TEST_ONNX_DIR})
 add_api_test(tf_parser test_tf_parser.cpp ${TEST_TF_DIR})
 # GPU-based tests

test/api/test_c_op_construct.c

+/*
+ * The MIT License (MIT)
+ *
+ * Copyright (c) 2015-2022 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
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ * THE SOFTWARE.
+ */
+#include <migraphx/migraphx.h>
+#include <string.h>
+
+void expect_equal(const char* x, const char* y)
+{
+    if(strcmp(x, y) != 0)
+        abort();
+}
+
+int main()
+{
+    char name[1024];
+    migraphx_operation_t op;
+    migraphx_operation_create(&op, "add", 0);
+    migraphx_operation_name(name, 1024, op);
+    migraphx_operation_destroy(op);
+    expect_equal(name, "add");
+}

tools/api/migraphx.h

@@ -25,6 +25,7 @@
 #define MIGRAPHX_GUARD_C_API_MIGRAPHX_H
 
 #include <stdlib.h>
+#include <stdbool.h>
 
 // Add new types here
 // clang-format off