Test the performance of permutation instances with different parameters

CMakeLists.txt

@@ -2,7 +2,7 @@
  #
  # MIT License
  #
- # Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ # Copyright (C) 2023-2025 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
@@ -65,6 +65,7 @@ set(CMAKE_CXX_FLAGS_MINSIZEREL "-O2 -DNDEBUG") # clang++ failed to build the pro
 if( CMAKE_PROJECT_NAME STREQUAL "hiptensor" )
   option( HIPTENSOR_BUILD_TESTS "Build hiptensor tests" ON )
   option( HIPTENSOR_BUILD_SAMPLES "Build hiptensor samples" ON )
+  option( HIPTENSOR_BUILD_TUNING "Build hiptensor tuning tools" ON )
   option( HIPTENSOR_BUILD_COMPRESSED_DBG "Enable compressed debug symbols" ON)
   option( HIPTENSOR_DATA_LAYOUT_COL_MAJOR "Set hiptensor data layout to column major" ON )
 endif()
@@ -166,6 +167,10 @@ if(HIPTENSOR_BUILD_SAMPLES)
   add_subdirectory(samples)
 endif()
 
+if(HIPTENSOR_BUILD_TUNING)
+	add_subdirectory(tuning)
+endif()
+
 # Versioning via rocm-cmake
 set ( VERSION_STRING "1.4.0" )
 rocm_setup_version( VERSION ${VERSION_STRING} )

library/src/permutation/hiptensor_permutation.cpp

@@ -233,8 +233,7 @@ hiptensorStatus_t hiptensorPermutation(const hiptensorHandle_t*           handle
                     return HIPTENSOR_STATUS_CK_ERROR;
                 }
 
-                int  n     = pSolution->problemDim();
-                auto flops = std::size_t(2) * n;
+                auto flops = std::size_t(2) * pSolution->problemSize();
                 auto bytes = pSolution->problemBytes();
 
                 hiptensor::PerfMetrics metrics = {

library/src/permutation/permutation_solution.cpp

@@ -2,7 +2,7 @@
  *
  * MIT License
  *
- * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ * Copyright (C) 2023-2025 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
@@ -33,6 +33,7 @@ namespace hiptensor
         std::unique_ptr<ck::tensor_operation::device::BaseOperator>&& deviceOp,
         std::unique_ptr<PermutationSolutionParams>&&                  params)
         : mDim(0)
+        , mSize(0)
         , mBytes(0)
         , mValid(false)
         , mDeviceOp(std::move(deviceOp))
@@ -42,6 +43,7 @@ namespace hiptensor
 
     PermutationSolution::PermutationSolution(PermutationSolution&& other)
         : mDim(other.mDim)
+        , mSize(other.mSize)
         , mBytes(other.mBytes)
         , mValid(other.mValid)
         , mDeviceOp(std::move(other.mDeviceOp))
@@ -57,6 +59,7 @@ namespace hiptensor
         {
             mDim = other.mDim;
 
+            mSize  = other.mSize;
             mBytes = other.mBytes;
             mValid = other.mValid;
 
@@ -142,6 +145,11 @@ namespace hiptensor
         return mDim;
     }
 
+    ck::index_t PermutationSolution::problemSize() const
+    {
+        return mSize;
+    }
+
     ck::index_t PermutationSolution::problemBytes() const
     {
         return mBytes;
@@ -167,6 +175,7 @@ namespace hiptensor
     void PermutationSolution::resetArgs()
     {
         mDim   = 0;
+        mSize  = 0;
         mBytes = 0;
 
         mInvokerArgPtr.reset(nullptr);

library/src/permutation/permutation_solution.hpp

@@ -2,7 +2,7 @@
  *
  * MIT License
  *
- * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ * Copyright (C) 2023-2025 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
@@ -105,6 +105,9 @@ namespace hiptensor
         // Problem dimension
         ck::index_t problemDim() const;
 
+        // Problem size
+        ck::index_t problemSize() const;
+
         // Byte count
         ck::index_t problemBytes() const;
 
@@ -120,6 +123,7 @@ namespace hiptensor
     protected:
         // Derived runtime arguments
         ck::index_t mDim;
+        ck::index_t mSize;
         ck::index_t mBytes;
         bool        mValid;
         uint32_t    mThreadDim;

library/src/permutation/permutation_solution_impl.hpp

@@ -2,7 +2,7 @@
  *
  * MIT License
  *
- * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ * Copyright (C) 2023-2025 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
@@ -150,9 +150,13 @@ namespace hiptensor
             // Fill problem metrics
             Base::mDim = Traits::NDim;
 
+            // Size count
+            Base::mSize
+                = std::accumulate(abLengths.cbegin(), abLengths.cend(), 1, std::multiplies{});
+
             // Byte count
-            Base::mBytes = sizeof(typename Traits::InDataT) * Base::mDim
-                           + sizeof(typename Traits::OutDataT) * Base::mDim;
+            Base::mBytes = (sizeof(typename Traits::InDataT) + sizeof(typename Traits::OutDataT))
+                           * Base::mSize;
 
             // Arg test
             Base::mValid = deviceOp->IsSupportedArgument(Base::mInvokerArgPtr.get());

samples/02_permutation/permutation.cpp

@@ -2,7 +2,7 @@
  *
  * MIT License
  *
- * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ * Copyright (C) 2023-2025 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
@@ -34,6 +34,7 @@
 
 #include <hiptensor/hiptensor.hpp>
 #include <hiptensor/internal/hiptensor_utility.hpp>
+#include <hiptensor_options.hpp>
 
 #include "common.hpp"
 
@@ -57,16 +58,15 @@ int main()
       B_{w, h, c, n} = 1.0 *  \textsl{IDENTITY}(A_{c, n, h, w})
      **********************/
 
-    std::vector<int> modeA{'w', 'h', 'c', 'n'};
-    std::vector<int> modeC{'c', 'n', 'h', 'w'};
+    std::vector<int> modeA{'w', 'h', 'c'};
+    std::vector<int> modeC{'c', 'w', 'h'};
     int              nmodeA = modeA.size();
     int              nmodeC = modeC.size();
 
     std::unordered_map<int, int64_t> extent;
-    extent['h'] = 32;
-    extent['w'] = 33;
-    extent['c'] = 34;
-    extent['n'] = 35;
+    extent['h'] = 512;
+    extent['w'] = 512;
+    extent['c'] = 512;
 
     std::vector<int64_t> extentA;
     for(auto mode : modeA)
@@ -117,6 +117,11 @@ int main()
     CHECK_HIPTENSOR_ERROR(hiptensorInitTensorDescriptor(
         handle, &descC, nmodeC, extentC.data(), NULL /* stride */, typeC, HIPTENSOR_OP_IDENTITY));
 
+    using hiptensor::HiptensorOptions;
+    auto& options = HiptensorOptions::instance();
+    options->setColdRuns(5);
+    options->setHotRuns(50);
+    hiptensorLoggerSetMask(0x1F);
     const floatTypeCompute one = 1.0f;
     CHECK_HIPTENSOR_ERROR(hiptensorPermutation(handle,
                                                &one,

samples/CMakeLists.txt

@@ -2,7 +2,7 @@
  #
  # MIT License
  #
- # Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ # Copyright (C) 2023-2025 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
@@ -44,6 +44,7 @@ if( CMAKE_PROJECT_NAME STREQUAL "hiptensor" )
         target_include_directories(${BINARY_NAME} PRIVATE
                                 ${CMAKE_CURRENT_SOURCE_DIR}
                                 ${PROJECT_SOURCE_DIR}/samples
+                                ${PROJECT_SOURCE_DIR}/library/src/include/
                                 ${PROJECT_SOURCE_DIR}/library/include)
 
         # Build this sample under custom target

tuning/CMakeLists.txt

@@ -0,0 +1,66 @@
+###############################################################################
+ #
+ # MIT License
+ #
+ # Copyright (C) 2023-2025 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.
+ #
+ ###############################################################################
+
+ # find_package( hip REQUIRED )
+find_package( composable_kernel 1.0.0 REQUIRED PATHS $ENV{CK_DIR}/lib/cmake /opt/rocm /opt/rocm/ck COMPONENTS device_contraction_operations device_reduction_operations device_other_operations)
+rocm_package_add_dependencies("composable_kernel >= 1.0.0" COMPONENT tests)
+
+set(THREADS_PREFER_PTHREAD_FLAG ON)
+find_package(Threads REQUIRED)
+
+# Create test executables and deploy
+function(add_hiptensor_component COMPONENT_NAME COMPONENT_SOURCES)
+
+    message( STATUS "adding hiptensor component: ${COMPONENT_NAME}")
+
+    # Ensure that all the sources are captured
+    list(APPEND COMPONENT_SOURCES ${ARGN})
+
+    # Build components as object files.
+    # Make sure they have -fPIC and that they inherit the hip::device environment
+    # so that they build for all archs in AMDGPU_TARGETS
+    add_library(${COMPONENT_NAME} OBJECT ${COMPONENT_SOURCES})
+    target_compile_options(${COMPONENT_NAME} PRIVATE ${CLANG_DRIVER_MODE})
+    target_link_options(${COMPONENT_NAME} PRIVATE ${CLANG_DRIVER_MODE})
+    set_target_properties(${COMPONENT_NAME} PROPERTIES POSITION_INDEPENDENT_CODE ON)
+    target_link_libraries(${COMPONENT_NAME} PRIVATE hip::device)
+    if(HIPTENSOR_BUILD_COMPRESSED_DBG)
+        target_compile_options(${COMPONENT_NAME} PRIVATE
+            "$<$<CONFIG:Debug>:-gz>"
+            "$<$<CONFIG:RelWithDebInfo>:-gz>"
+            )
+    endif()
+
+endfunction()
+
+
+# Generates hiptensor_contraction and hiptensor_contraction_instances
+# add_subdirectory(contraction)
+# Generates hiptensor_permutation and hiptensor_permutation_instances
+add_subdirectory(permutation)
+# Generates hiptensor_reduction and hiptensor_reduction_instances
+# add_subdirectory(reduction)
+

tuning/permutation/CMakeLists.txt

@@ -0,0 +1,66 @@
+###############################################################################
+#
+# MIT License
+#
+# Copyright (C) 2023-2025 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.
+#
+###############################################################################
+
+# Don't link to consume ck instances, however will customize and create our own.
+# Make the ck includes visible so we can build instances.
+get_target_property(composable_kernel_INCLUDES composable_kernel::device_other_operations INTERFACE_INCLUDE_DIRECTORIES)
+
+
+function(add_permutation_tuning RANK)
+
+	set(HIPTENSOR_PERMUTATION_TUNING "permutation_tuning_${RANK}")
+
+	set(HIPTENSOR_PERMUTATION_TUNING_SOURCES
+		${CMAKE_CURRENT_SOURCE_DIR}/permutation_instance_${RANK}.cpp
+		${CMAKE_CURRENT_SOURCE_DIR}/permutation_instance_F16_${RANK}_256_64_64_4_4.cpp
+		${CMAKE_CURRENT_SOURCE_DIR}/permutation_instance_F32_${RANK}_256_64_64_4_4.cpp
+		${CMAKE_CURRENT_SOURCE_DIR}/permutation_instance_F16_${RANK}_256_64_64_16_16.cpp
+		${CMAKE_CURRENT_SOURCE_DIR}/permutation_instance_F32_${RANK}_256_64_64_16_16.cpp
+		${CMAKE_CURRENT_SOURCE_DIR}/permutation_instance_F16_${RANK}_256_128_128_8_8.cpp
+		${CMAKE_CURRENT_SOURCE_DIR}/permutation_instance_F32_${RANK}_256_128_128_8_8.cpp
+		${CMAKE_CURRENT_SOURCE_DIR}/permutation_instance_F16_${RANK}_256_128_128_16_16.cpp
+		${CMAKE_CURRENT_SOURCE_DIR}/permutation_instance_F32_${RANK}_256_128_128_16_16.cpp
+		${CMAKE_CURRENT_SOURCE_DIR}/permutation_instance_F16_${RANK}_miscellaneous.cpp
+		${CMAKE_CURRENT_SOURCE_DIR}/permutation_instance_F32_${RANK}_miscellaneous.cpp
+		${CMAKE_CURRENT_SOURCE_DIR}/permutation_tuning_${RANK}.cpp
+		)
+
+	add_executable(${HIPTENSOR_PERMUTATION_TUNING} ${HIPTENSOR_PERMUTATION_TUNING_SOURCES})
+
+	set_target_properties(${HIPTENSOR_PERMUTATION_TUNING} PROPERTIES POSITION_INDEPENDENT_CODE ON)
+	target_compile_options(${HIPTENSOR_PERMUTATION_TUNING} PRIVATE ${CMAKE_CXX_FLAGS} ${CLANG_DRIVER_MODE})
+	target_link_options(${HIPTENSOR_PERMUTATION_TUNING} PRIVATE ${CLANG_DRIVER_MODE})
+	target_include_directories(${HIPTENSOR_PERMUTATION_TUNING} PRIVATE ${composable_kernel_INCLUDES})
+	target_include_directories(${HIPTENSOR_PERMUTATION_TUNING} PRIVATE "${CMAKE_SOURCE_DIR}/library/src/include/" )
+	target_link_libraries(${HIPTENSOR_PERMUTATION_TUNING} PRIVATE hip::device hip::host)
+
+endfunction()
+
+add_permutation_tuning(2)
+add_permutation_tuning(3)
+add_permutation_tuning(4)
+add_permutation_tuning(5)
+add_permutation_tuning(6)