From 50fc7ddeea72bfdfed951099a0c85194c048d88c Mon Sep 17 00:00:00 2001
From: Jianfeng yan <jfyan008@gmail.com>
Date: Fri, 11 Mar 2022 06:24:02 +0000
Subject: [PATCH 1/2] added more static_assert

---
 .../gpu/thread/threadwise_tensor_slice_transfer.hpp   | 11 ++++++++++-
 include/ck/utility/sequence.hpp                       |  6 ++++++
 include/ck/utility/tensor_space_filling_curve.hpp     |  4 ++++
 3 files changed, 20 insertions(+), 1 deletion(-)

diff --git a/include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer.hpp b/include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer.hpp
index 524da47e245..48e3f4afd68 100644
--- a/include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer.hpp
+++ b/include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer.hpp
@@ -79,6 +79,8 @@ struct ThreadwiseTensorSliceTransfer_v1r3
     {
         static_assert(SrcDesc::IsKnownAtCompileTime(),
                       "wrong! SrcDesc need to known at compile-time");
+        static_assert(SliceLengths::At(Number<DstVectorDim>{}) % DstScalarPerVector == 0,
+                      "wrong! Not divisible");
     }
 
     __device__ void SetDstSliceOrigin(const DstDesc& dst_desc, const Index& dst_slice_origin_idx)
@@ -250,6 +252,8 @@ struct ThreadwiseTensorSliceTransfer_v2
     {
         static_assert(DstDesc::IsKnownAtCompileTime(),
                       "wrong! SrcDesc need to known at compile-time");
+        static_assert(SliceLengths::At(Number<SrcVectorDim>{}) % SrcScalarPerVector == 0,
+                      "wrong! Not divisible");
     }
 
     __device__ void SetSrcSliceOrigin(const SrcDesc& src_desc, const Index& src_slice_origin_idx)
@@ -439,6 +443,10 @@ struct ThreadwiseTensorSliceTransfer_v3
         : src_coord_(make_tensor_coordinate(src_desc, src_slice_origin)),
           dst_coord_(make_tensor_coordinate(dst_desc, dst_slice_origin))
     {
+        static_assert(SliceLengths::At(Number<SrcVectorDim>{}) % SrcScalarPerVector == 0,
+                      "wrong! Not divisible");
+        static_assert(SliceLengths::At(Number<DstVectorDim>{}) % DstScalarPerVector == 0,
+                      "wrong! Not divisible");
     }
 
     __device__ void SetSrcSliceOrigin(const SrcDesc& src_desc, const Index& src_slice_origin_idx)
@@ -1016,7 +1024,8 @@ struct ThreadwiseTensorSliceTransfer_v4
         static_assert(SrcDesc::IsKnownAtCompileTime() && DstDesc::IsKnownAtCompileTime(),
                       "wrong! SrcDesc and DstDesc need to known at compile-time");
 
-        static_assert(SliceLengths::At(Number<SrcVectorDim>{}) % SrcScalarPerVector == 0, "wrong!");
+        static_assert(SliceLengths::At(Number<SrcVectorDim>{}) % SrcScalarPerVector == 0,
+                      "wrong! Not divisible");
     }
 
     template <typename SrcRefToOriginDisplacement,
diff --git a/include/ck/utility/sequence.hpp b/include/ck/utility/sequence.hpp
index b35999d56ff..85eb3a53d38 100644
--- a/include/ck/utility/sequence.hpp
+++ b/include/ck/utility/sequence.hpp
@@ -606,6 +606,12 @@ struct sequence_map_inverse
                                            SeqMap::Size()>::type;
 };
 
+template <index_t... Xs, index_t... Ys>
+__host__ __device__ constexpr bool operator==(Sequence<Xs...>, Sequence<Ys...>)
+{
+    return (true && ... && (Xs == Ys));
+}
+
 template <index_t... Xs, index_t... Ys>
 __host__ __device__ constexpr auto operator+(Sequence<Xs...>, Sequence<Ys...>)
 {
diff --git a/include/ck/utility/tensor_space_filling_curve.hpp b/include/ck/utility/tensor_space_filling_curve.hpp
index c5cbe461f0b..62b68559bf0 100644
--- a/include/ck/utility/tensor_space_filling_curve.hpp
+++ b/include/ck/utility/tensor_space_filling_curve.hpp
@@ -37,6 +37,10 @@ struct SpaceFillingCurve
 
     __host__ __device__ static constexpr index_t GetNumOfAccess()
     {
+        static_assert(TensorLengths::Size() == ScalarsPerAccess::Size());
+        static_assert(TensorLengths{} % ScalarsPerAccess{} ==
+                      typename uniform_sequence_gen<TensorLengths::Size(), 0>::type{});
+
         return reduce_on_sequence(TensorLengths{}, math::multiplies{}, Number<1>{}) /
                ScalarPerVector;
     }

From 43c8b6d03cddbbb5fbe370f1e6ae4177d608bf50 Mon Sep 17 00:00:00 2001
From: Jianfeng yan <jfyan008@gmail.com>
Date: Fri, 11 Mar 2022 17:52:16 +0000
Subject: [PATCH 2/2] minor change

---
 include/ck/utility/sequence.hpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/include/ck/utility/sequence.hpp b/include/ck/utility/sequence.hpp
index 85eb3a53d38..c2adfc5063f 100644
--- a/include/ck/utility/sequence.hpp
+++ b/include/ck/utility/sequence.hpp
@@ -609,7 +609,7 @@ struct sequence_map_inverse
 template <index_t... Xs, index_t... Ys>
 __host__ __device__ constexpr bool operator==(Sequence<Xs...>, Sequence<Ys...>)
 {
-    return (true && ... && (Xs == Ys));
+    return ((Xs == Ys) && ...);
 }
 
 template <index_t... Xs, index_t... Ys>