Reference kernel for 3D convolution for non-packed tensors

src/CMakeLists.txt

@@ -388,6 +388,7 @@ if( MIOPEN_BACKEND MATCHES "OpenCL" OR MIOPEN_BACKEND STREQUAL "HIPOC" OR MIOPEN
         kernels/workaround_issue_1431.hpp
         kernels/hip_f8_impl.hpp
         kernels/hip_float8.hpp
+        kernels/stride_array.hpp
         )
 
     set(MIOPEN_KERNELS

src/include/miopen/hipoc_kernel.hpp

@@ -26,14 +26,15 @@
 #ifndef GUARD_MIOPEN_HIPOC_KERNEL_HPP
 #define GUARD_MIOPEN_HIPOC_KERNEL_HPP
 
-#include <array>
-#include <cassert>
 #include <miopen/errors.hpp>
 #include <miopen/hipoc_program.hpp>
 #include <miopen/stringutils.hpp>
 #include <miopen/op_kernel_args.hpp>
+
+#include <array>
+#include <cassert>
+#include <cstring>
 #include <vector>
-#include <memory.h>
 
 namespace miopen {
 
@@ -47,36 +48,20 @@ inline HipEventPtr make_hip_event()
 
 #if 1 // Keep around other storage techinques -- @pfultz2 27.03.2017
 
-#if 0 // Keep around other storage techinques -- @pfultz2 27.03.2017
-template <class T, class U>
-struct KernelArgsPair
-{
-    static const int alignment    = sizeof(U);
-    static const int padding      = (alignment - sizeof(T) % alignment) % alignment;
-    static const int second_index = sizeof(T) + padding;
-    KernelArgsPair(T x, U y)
-    {
-        new(buffer) T(x); // NOLINT (clang-analyzer-cplusplus.PlacementNew)
-        new(buffer + second_index) U(y);
-    }
-    alignas(U) char buffer[second_index + sizeof(U)] = {};
-};
-#else
 template <class T, class U>
 struct KernelArgsPair
 {
-    static const int alignment = alignof(U);
-    static const int padding   = (alignment - (sizeof(T) % alignment)) % alignment;
-    static_assert(padding >= 0, "padding cannot be negative");
-    static const int second_index = sizeof(T) + padding;
+    constexpr static const auto alignU       = alignof(U);
+    constexpr static const auto padding      = (alignU - (sizeof(T) % alignU)) % alignU;
+    constexpr static const auto second_index = sizeof(T) + padding;
     KernelArgsPair(T x, U y)
     {
         new(buffer) T(x); // NOLINT (clang-analyzer-cplusplus.PlacementNew)
         new(buffer + second_index) U(y);
     }
+
     alignas(U) char buffer[second_index + sizeof(U)] = {};
 };
-#endif
 
 template <class... Ts>
 struct KernelArgsPack;

src/include/miopen/solver/conv_direct_naive_conv.hpp

@@ -27,12 +27,13 @@
 
 #include <miopen/execution_context.hpp>
 #include <miopen/problem_description.hpp>
+#include "miopen/../../kernels/stride_array.hpp"
 
 #include <array>
 #include <algorithm>
+#include <cassert>
 #include <string>
 #include <vector>
-#include <cassert>
 
 namespace miopen {
 
@@ -65,54 +66,6 @@ void printTensorStrides(const TensorDescriptor& inDesc,
                         const TensorDescriptor& wDesc,
                         const TensorDescriptor& outDesc);
 
-// TODO(Amber): Uncomment when hip RTC accepts std::array
-// using StrideIndexType = int;
-// using Strides3D       = std::array<StrideIndexType, 3>;
-// using Strides4D       = std::array<StrideIndexType, 4>;
-// using Strides5D       = std::array<StrideIndexType, 5>;
-// using Strides6D       = std::array<StrideIndexType, 6>;
-#if 1
-template <typename T, unsigned N>
-class MyArray
-{
-    T data_[N] = {};
-
-public:
-    constexpr static const unsigned SIZE = N;
-    __host__ __device__ constexpr unsigned size() const { return N; }
-
-    __host__ __device__ const T& operator[](unsigned i) const { return data_[i]; }
-
-    __host__ T& operator[](unsigned i) { return data_[i]; }
-
-    __host__ __device__ MyArray()                   = default;
-    __host__ __device__ MyArray(const MyArray&)     = default;
-    __host__ __device__ MyArray(MyArray&&) noexcept = default;
-    __host__ __device__ MyArray& operator=(const MyArray&) = default;
-    __host__ __device__ MyArray& operator=(MyArray&&) noexcept = default;
-    __host__ __device__ ~MyArray()                             = default;
-};
-
-using StrideIndexType = size_t;
-using Strides5D       = MyArray<StrideIndexType, 5u>;
-using Strides6D       = MyArray<StrideIndexType, 6u>;
-
-#else
-
-extern "C" typedef int StrideIndexType;
-
-extern "C" typedef struct
-{
-    StrideIndexType v[5];
-} Strides5D;
-
-extern "C" typedef struct
-{
-    StrideIndexType v[6];
-} Strides6D;
-
-#endif
-
 namespace internal {
 template <unsigned N>
 struct ChooseStride
@@ -189,27 +142,5 @@ V SplitWeiStrideKtoGK(int k_per_group, const V& wei_strides)
     return ret;
 }
 
-template <typename StrideArray>
-void printStrideArray(const char* name, const StrideArray& sarr)
-{
-    printf("%s = [", name);
-    for(unsigned i = 0; i < StrideArray::SIZE; ++i)
-    {
-        printf("%d,", sarr[i]);
-    }
-    printf("]\n");
-}
-
-template <typename StrideArray>
-void printStrideArrays(const StrideArray& in_strides,
-                       const StrideArray& wei_strides,
-                       const StrideArray& out_strides)
-{
-
-    printStrideArray("in_strides", in_strides);
-    printStrideArray("wei_strides", wei_strides);
-    printStrideArray("out_strides", out_strides);
-}
-
 } // namespace solver
 } // namespace miopen