fix comments, remove redundant code and fix formatting issues

rapidsai · rapids-bot · Nov 16, 2022 · Oct 20, 2022 · Oct 21, 2022 · Oct 21, 2022
commit 7786fcb882bda7a9b243a57de5544f1a09b349de
@@ -26,12 +26,14 @@ namespace detail {
 
 template <typename DataT, typename AccT>
 struct CosineOp {
-  __device__ __host__ CosineOp() noexcept {}
-  __device__ __host__ AccT operator()(DataT& aNorm, const DataT& bNorm, DataT& accVal) const noexcept
+  __device__  CosineOp() noexcept {}
+  __device__  AccT operator()(DataT& aNorm,
+                                      const DataT& bNorm,
+                                      DataT& accVal) const noexcept
   {
     return static_cast<AccT>(1.0) - (AccT)(accVal / (aNorm * bNorm));
   }
-  __device__ __host__ AccT operator()(DataT aData) const noexcept { return aData; }
+  __device__  AccT operator()(DataT aData) const noexcept { return aData; }
 };
 
 /**

@@ -616,15 +616,15 @@ void distance(const InType* x,
  *  worksize, the number of bytes of workspace required
  */
 
-// Default final op functor which facilitates elementwise operation on 
+// Default final op functor which facilitates elementwise operation on
 // final distance value if any.
 template <typename AccType, typename OutType, typename Index>
 struct default_fin_op {
   __host__ __device__ default_fin_op() noexcept {};
   // functor signature.
   __host__ __device__ OutType operator()(AccType d_val, Index g_d_idx) const noexcept
   {
-    return d_val; 
+    return d_val;
   }
 };
 

@@ -28,15 +28,17 @@ template <typename DataT, typename AccT>
 struct L2ExpandedOp {
   bool sqrt;
 
-  __device__ __host__ L2ExpandedOp() noexcept : sqrt(false) {}
-  __device__ __host__ L2ExpandedOp(bool isSqrt) noexcept : sqrt(isSqrt) {}
-  __device__ __host__ AccT operator()(DataT& aNorm, const DataT& bNorm, DataT& accVal) const noexcept
+  __device__  L2ExpandedOp() noexcept : sqrt(false) {}
+  __device__  L2ExpandedOp(bool isSqrt) noexcept : sqrt(isSqrt) {}
+  __device__  AccT operator()(DataT& aNorm,
+                                      const DataT& bNorm,
+                                      DataT& accVal) const noexcept
   {
     AccT outVal = aNorm + bNorm - DataT(2.0) * accVal;
     return sqrt ? raft::mySqrt(outVal) : outVal;
   }
 
-  __device__ __host__ AccT operator()(DataT aData) const noexcept { return aData; }
+  __device__  AccT operator()(DataT aData) const noexcept { return aData; }
 };
 
 /**

@@ -70,6 +70,9 @@ void cutlassDistanceKernel(const DataT* x,
                            DistanceFn dist_op,
                            cudaStream_t stream)
 {
+  static_assert(!(std::is_same<OutT, bool>::value),
+                "OutType bool is not supported use uint8_t instead");
+
   using EpilogueOutputOp =
     cutlass::epilogue::thread::PairwiseDistanceEpilogueElementwise<DataT,  // ElementC_
                                                                    AccT,   // ElementAccumulator_
@@ -95,8 +98,7 @@ void cutlassDistanceKernel(const DataT* x,
   constexpr int Alignment = VecLen;
 
   // default initialize problem size with row major inputs
-  auto problem_size =
-    cutlass::gemm::GemmCoord(static_cast<int>(n), static_cast<int>(m), static_cast<int>(k));
+  auto problem_size = cutlass::gemm::GemmCoord(n, m, k);
 
   using cutlassDistKernel =
     typename cutlass::gemm::kernel::PairwiseDistanceGemm<DataT,
@@ -117,12 +119,11 @@ void cutlassDistanceKernel(const DataT* x,
     gemm_lda = ldb;
     gemm_ldb = lda;
   } else {
-    problem_size =
-      cutlass::gemm::GemmCoord(static_cast<int>(m), static_cast<int>(n), static_cast<int>(k));
-    a        = x;
-    b        = y;
-    gemm_lda = lda;
-    gemm_ldb = ldb;
+    problem_size = cutlass::gemm::GemmCoord(m, n, k);
+    a            = x;
+    b            = y;
+    gemm_lda     = lda;
+    gemm_ldb     = ldb;
   }
 
   typename cutlassDist::Arguments arguments{

@@ -1,33 +1,19 @@
-/***************************************************************************************************
- * Copyright (c) 2017 - 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
- * SPDX-License-Identifier: BSD-3-Clause
+/*
+ * Copyright (c) 2018-2022, NVIDIA CORPORATION.
  *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions are met:
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
  *
- * 1. Redistributions of source code must retain the above copyright notice, this
- * list of conditions and the following disclaimer.
+ *     http://www.apache.org/licenses/LICENSE-2.0
  *
- * 2. Redistributions in binary form must reproduce the above copyright notice,
- * this list of conditions and the following disclaimer in the documentation
- * and/or other materials provided with the distribution.
- *
- * 3. Neither the name of the copyright holder nor the names of its
- * contributors may be used to endorse or promote products derived from
- * this software without specific prior written permission.
- *
- * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
- * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
- * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
- * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
- * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
- * DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
- * SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
- * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
- * OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
- * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- *
- **************************************************************************************************/
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
 /*! \file
   \brief Epilogue for threadblock scoped GEMMs using Tensor Ops.
 

@@ -1,3 +1,19 @@
+/*
+ * Copyright (c) 2018-2022, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
 //
 /*! \file
   \brief Functor performing distance operations used by epilogues of pairwise distance
@@ -106,18 +122,7 @@ class PairwiseDistanceEpilogueElementwise {
 
   /// Functionally required for serial reduction in the epilogue
   CUTLASS_HOST_DEVICE
-  void set_k_partition(int k_partition, int k_partition_count)
-  {
-#if 0
-    if (k_partition) {
-      beta_ = ElementCompute(1);
-    }
-
-    if (k_partition != k_partition_count - 1) {
-      skip_elementwise_ = true;
-    }
-#endif
-  }
+  void set_k_partition(int k_partition, int k_partition_count) {}
 
   /// Applies the operation when is_source_needed() is true
   CUTLASS_HOST_DEVICE

@@ -1,3 +1,19 @@
+/*
+ * Copyright (c) 2018-2022, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
 #pragma once
 
 #include "cutlass/cutlass.h"
@@ -7,7 +23,6 @@
 #include "cutlass/layout/matrix.h"
 #include "cutlass/layout/tensor.h"
 
-//#include "./epilogue_with_bcast_threadblock.h"
 #include "pairwise_distance_epilogue.h"
 
 /////////////////////////////////////////////////////////////////////////////////////////////////
@@ -44,7 +59,7 @@ struct PairwiseDistanceGemm {
 
   /// Threadblock-level tile size (concept: GemmShape)
   using ThreadblockShape =
-    cutlass::gemm::GemmShape<128, 128, 16>;  // <- threadblock tile M = 128, N = 64, K = 16
+    cutlass::gemm::GemmShape<128, 128, 16>;  // <- threadblock tile M = 128, N = 128, K = 16
   /// Warp-level tile size (concept: GemmShape)
   // This code section describes tile size a warp will compute
   using WarpShape = cutlass::gemm::GemmShape<64, 64, 16>;  // <- warp tile M = 64, N = 64, K = 16
@@ -145,10 +160,10 @@ struct PairwiseDistanceGemm<double,
   // using Transform = cutlass::ComplexTransform::kNone;
   // Threadblock-level tile size (concept: GemmShape)
   using ThreadblockShape =
-    cutlass::gemm::GemmShape<64, 64, 16>;  // <- threadblock tile M = 128, N = 64, K = 16
+    cutlass::gemm::GemmShape<64, 64, 16>;  // <- threadblock tile M = 64, N = 64, K = 16
   /// Warp-level tile size (concept: GemmShape)
   // This code section describes tile size a warp will compute
-  using WarpShape = cutlass::gemm::GemmShape<32, 32, 16>;  // <- warp tile M = 64, N = 64, K = 16
+  using WarpShape = cutlass::gemm::GemmShape<32, 32, 16>;  // <- warp tile M = 32, N = 32, K = 16
   /// Warp-level tile size (concept: GemmShape)
   // This code section describes the size of MMA op
   using InstructionShape = cutlass::gemm::GemmShape<8, 8, 4>;

@@ -1,33 +1,19 @@
-/***************************************************************************************************
- * Copyright (c) 2017 - 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
- * SPDX-License-Identifier: BSD-3-Clause
+/*
+ * Copyright (c) 2018-2022, NVIDIA CORPORATION.
  *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions are met:
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
  *
- * 1. Redistributions of source code must retain the above copyright notice, this
- * list of conditions and the following disclaimer.
+ *     http://www.apache.org/licenses/LICENSE-2.0
  *
- * 2. Redistributions in binary form must reproduce the above copyright notice,
- * this list of conditions and the following disclaimer in the documentation
- * and/or other materials provided with the distribution.
- *
- * 3. Neither the name of the copyright holder nor the names of its
- * contributors may be used to endorse or promote products derived from
- * this software without specific prior written permission.
- *
- * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
- * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
- * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
- * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
- * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
- * DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
- * SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
- * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
- * OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
- * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- *
- **************************************************************************************************/
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
 /*! \file
   \brief Epilogue for threadblock scoped GEMMs using Tensor Ops.
 
@@ -78,7 +64,6 @@ class PredicatedTileIteratorNormVec {
 
   using Element = Element_;
 
-  // using Layout = layout::RowMajor;
   using Layout         = Layout_;
   using TensorRef      = TensorRef<Element, Layout>;
   using ConstTensorRef = typename TensorRef::ConstTensorRef;

diff --git a/cpp/include/raft/distance/distance.cuh b/cpp/include/raft/distance/distance.cuh
@@ -315,7 +315,6 @@ void pairwise_distance(const raft::handle_t& handle,
       detail::pairwise_distance_impl<Type, Index_, raft::distance::DistanceType::CosineExpanded>(
         x, y, dist, m, n, k, workspace, handle.get_stream(), isRowMajor);
       break;
-#if 0
     case raft::distance::DistanceType::L1:
       detail::pairwise_distance_impl<Type, Index_, raft::distance::DistanceType::L1>(
         x, y, dist, m, n, k, workspace, handle.get_stream(), isRowMajor);
@@ -365,7 +364,6 @@ void pairwise_distance(const raft::handle_t& handle,
         pairwise_distance_impl<Type, Index_, raft::distance::DistanceType::CorrelationExpanded>(
           x, y, dist, m, n, k, workspace, handle.get_stream(), isRowMajor);
       break;
-#endif
     default: THROW("Unknown or unsupported distance metric '%d'!", (int)metric);
   };
 }