Imprve indexing and syncthreads

src/targets/gpu/compile_hip.cpp

@@ -43,6 +43,7 @@ inline namespace MIGRAPHX_INLINE_NS {
 namespace gpu {
 
 MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_GPU_DEBUG);
+MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_GPU_DEBUG_SYM);
 MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_GPU_OPTIMIZE);
 MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_GPU_DUMP_ASM);
 MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_GPU_DUMP_SRC);
@@ -227,6 +228,8 @@ compile_hip_src(const std::vector<src_file>& srcs, std::string params, const std
     if(params.find("-std=") == std::string::npos)
         params += " --std=c++17";
     params += " -fno-gpu-rdc";
+    if(enabled(MIGRAPHX_GPU_DEBUG_SYM{}))
+        params += " -g";
     params += " -c";
     if(is_hcc_compiler())
     {

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

@@ -76,6 +76,7 @@ namespace migraphx {
 template <class T, index_int N>
 struct array
 {
+    using type = T;
     T d[N];
     constexpr T& operator[](index_int i)
     {
@@ -208,6 +209,15 @@ struct integral_const_array : array<T, sizeof...(Xs)>
     MIGRAPHX_DEVICE_CONSTEXPR integral_const_array() : base_array({Xs...}) {}
 };
 
+template<class F>
+constexpr auto return_const_array(F f)
+{
+    constexpr const auto a = f();
+    return sequence(a.size(), [=](auto... is) {
+        return integral_const_array<typename decltype(a)::type, a[is]...>{};
+    });
+}
+
 template <class T, T... Xs, class F>
 constexpr auto transform(integral_const_array<T, Xs...>, F f)
 {

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

@@ -111,6 +111,12 @@ struct gens<1> : seq<0>
 {
 };
 
+template <class F, index_int... Ns>
+constexpr void repeat_c_impl(F f, seq<Ns...>)
+{
+    swallow{(f(std::integral_constant<std::size_t, Ns>{}), 0)...};
+}
+
 template <class F, index_int... Ns>
 constexpr auto sequence_c_impl(F&& f, seq<Ns...>)
 {
@@ -143,6 +149,18 @@ constexpr auto sequence(IntegerConstant ic, F&& f)
     return sequence_c<ic>(f);
 }
 
+template <std::size_t N, class F>
+constexpr void repeat_c(F f)
+{
+    detail::repeat_c_impl(f, detail::gens<N>{});
+}
+
+template <class IntegerConstant, class F>
+constexpr auto repeat(IntegerConstant ic, F&& f)
+{
+    return repeat_c<ic>(f);
+}
+
 template <class F, class G>
 constexpr auto by(F f, G g)
 {

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

@@ -54,8 +54,8 @@ struct index
     }
 #endif
 
-    template <class F>
-    __device__ void global_stride(index_int n, F f) const
+    template <class N, class F>
+    __device__ void global_stride(N n, F f) const
     {
         const auto stride = nglobal();
         for(index_int i = global; i < n; i += stride)
@@ -64,8 +64,8 @@ struct index
         }
     }
 
-    template <class F>
-    __device__ void local_stride(index_int n, F f) const
+    template <class N, class F>
+    __device__ void local_stride(N n, F f) const
     {
         const auto stride = nlocal();
         for(index_int i = local; i < n; i += stride)

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/kernels/include/migraphx/kernels/shape.hpp

@@ -74,6 +74,20 @@ struct shape
 
     constexpr auto standard() const { return packed() and not transposed(); }
 
+    static constexpr auto compute_standard_strides()
+    {
+        return return_const_array([] {
+            index_array result{};
+            index_int s = 1;
+            for(diff_int is = result.size() - 1; is >= 0; is--)
+            {
+                result[is] = s;
+                s *= Lens{}[is];
+            }
+            return result;
+        });
+    }
+
     constexpr index_int index(index_array x) const { return x.dot(strides); }
 
     constexpr index_int index(std::initializer_list<index_int> x) const
@@ -86,8 +100,7 @@ struct shape
 
     constexpr index_int index(index_int i) const
     {
-        if(this->standard())
-        {
+        if constexpr(shape{}.standard())
             MIGRAPHX_ASSERT(i == compute_index(i));
             return i;
         }
@@ -99,19 +112,21 @@ struct shape
 
     constexpr index_int compute_index(index_int i) const
     {
-        const auto rank  = this->lens.size();
-        index_int s      = 1;
+        constexpr auto sstrides = compute_standard_strides();
+        constexpr auto rank  = Lens{}.size();
         index_int result = 0;
-        for(index_int j = 0; j < rank; j++)
-        {
-            const index_int k      = rank - j - 1;
-            const index_int stride = this->strides[k];
-            const index_int len    = this->lens[k];
-            const index_int slen   = s * len;
-            const index_int idx    = (i % slen) / s;
-            result += stride * idx;
-            s = slen;
-        }
+        repeat(rank, [&](auto j) {
+            constexpr auto k      = rank - j - 1;
+            constexpr auto stride = Strides{}[k];
+            constexpr auto len    = Lens{}[k];
+            if constexpr(stride != 0)
+            {
+                constexpr auto s   = sstrides[k];
+                constexpr auto slen   = s * len;
+                auto idx    = (i % slen) / s;
+                result += stride * idx;
+            }
+        });
         return result;
     }