Add tuning tool for permutation rank3

tuning/permutation/CMakeLists.txt

@@ -28,17 +28,35 @@
 # Make the ck includes visible so we can build instances.
 get_target_property(composable_kernel_INCLUDES composable_kernel::device_other_operations INTERFACE_INCLUDE_DIRECTORIES)
 
-set(HIPTENSOR_PERMUTATION_TUNING "permutation_tuning")
 
-set(HIPTENSOR_PERMUTATION_TUNING_SOURCES
-   ${CMAKE_CURRENT_SOURCE_DIR}/permutation_tuning.cpp
-)
+function(add_permutation_tuning RANK)
 
-add_executable(${HIPTENSOR_PERMUTATION_TUNING} ${HIPTENSOR_PERMUTATION_TUNING_SOURCES})
+	set(HIPTENSOR_PERMUTATION_TUNING "permutation_tuning_${RANK}")
 
-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_link_libraries(${HIPTENSOR_PERMUTATION_TUNING} PRIVATE hip::device hip::host)
+	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(3)

tuning/permutation/permutation_instance_3.cpp

@@ -0,0 +1,84 @@
+#include "permutation_instance.hpp"
+
+namespace hiptensor
+{
+    namespace tuning
+    {
+
+        namespace permutation
+        {
+
+            void genInstances_F16_3_256_128_128_16_16(
+                std::vector<std::unique_ptr<
+                    DeviceElementwise<ck::Tuple<F16>, ck::Tuple<F16>, UnaryScaleSquare, 3>>>&
+                    instances);
+            void genInstances_F16_3_256_128_128_8_8(
+                std::vector<std::unique_ptr<
+                    DeviceElementwise<ck::Tuple<F16>, ck::Tuple<F16>, UnaryScaleSquare, 3>>>&
+                    instances);
+            void genInstances_F16_3_256_64_64_16_16(
+                std::vector<std::unique_ptr<
+                    DeviceElementwise<ck::Tuple<F16>, ck::Tuple<F16>, UnaryScaleSquare, 3>>>&
+                    instances);
+            void genInstances_F16_3_256_64_64_4_4(
+                std::vector<std::unique_ptr<
+                    DeviceElementwise<ck::Tuple<F16>, ck::Tuple<F16>, UnaryScaleSquare, 3>>>&
+                    instances);
+            void genInstances_F16_3_miscellaneous(
+                std::vector<std::unique_ptr<
+                    DeviceElementwise<ck::Tuple<F16>, ck::Tuple<F16>, UnaryScaleSquare, 3>>>&
+                    instances);
+            void genInstances_F32_3_256_128_128_16_16(
+                std::vector<std::unique_ptr<
+                    DeviceElementwise<ck::Tuple<F32>, ck::Tuple<F32>, UnaryScaleSquare, 3>>>&
+                    instances);
+            void genInstances_F32_3_256_128_128_8_8(
+                std::vector<std::unique_ptr<
+                    DeviceElementwise<ck::Tuple<F32>, ck::Tuple<F32>, UnaryScaleSquare, 3>>>&
+                    instances);
+            void genInstances_F32_3_256_64_64_16_16(
+                std::vector<std::unique_ptr<
+                    DeviceElementwise<ck::Tuple<F32>, ck::Tuple<F32>, UnaryScaleSquare, 3>>>&
+                    instances);
+            void genInstances_F32_3_256_64_64_4_4(
+                std::vector<std::unique_ptr<
+                    DeviceElementwise<ck::Tuple<F32>, ck::Tuple<F32>, UnaryScaleSquare, 3>>>&
+                    instances);
+            void genInstances_F32_3_miscellaneous(
+                std::vector<std::unique_ptr<
+                    DeviceElementwise<ck::Tuple<F32>, ck::Tuple<F32>, UnaryScaleSquare, 3>>>&
+                    instances);
+
+            std::vector<std::unique_ptr<
+                DeviceElementwise<ck::Tuple<F16>, ck::Tuple<F16>, UnaryScaleSquare, 3>>>
+                genInstances_F16_3()
+            {
+                std::vector<std::unique_ptr<
+                    DeviceElementwise<ck::Tuple<F16>, ck::Tuple<F16>, UnaryScaleSquare, 3>>>
+                    instances;
+                genInstances_F16_3_256_64_64_4_4(instances);
+                genInstances_F16_3_256_64_64_16_16(instances);
+                genInstances_F16_3_256_128_128_8_8(instances);
+                genInstances_F16_3_256_128_128_16_16(instances);
+                genInstances_F16_3_miscellaneous(instances);
+                return instances;
+            }
+
+            std::vector<std::unique_ptr<
+                DeviceElementwise<ck::Tuple<F32>, ck::Tuple<F32>, UnaryScaleSquare, 3>>>
+                genInstances_F32_3()
+            {
+                std::vector<std::unique_ptr<
+                    DeviceElementwise<ck::Tuple<F32>, ck::Tuple<F32>, UnaryScaleSquare, 3>>>
+                    instances;
+                genInstances_F32_3_256_64_64_4_4(instances);
+                genInstances_F32_3_256_64_64_16_16(instances);
+                genInstances_F32_3_256_128_128_8_8(instances);
+                genInstances_F32_3_256_128_128_16_16(instances);
+                genInstances_F32_3_miscellaneous(instances);
+                return instances;
+            }
+
+        }
+    }
+}

tuning/permutation/permutation_instance_F16_3_256_128_128_16_16.cpp

@@ -0,0 +1,13 @@
+#include "permutation_instance.hpp"
+
+namespace hiptensor::tuning::permutation
+{
+
+    void genInstances_F16_3_256_128_128_16_16(
+        std::vector<std::unique_ptr<
+            DeviceElementwise<ck::Tuple<F16>, ck::Tuple<F16>, UnaryScaleSquare, 3>>>& instances)
+    {
+        genInstances_256_128_128_16_16<F16, 3>(instances);
+    }
+
+}

tuning/permutation/permutation_instance_F16_3_256_128_128_32_32.cpp

@@ -0,0 +1,13 @@
+#include "permutation_instance.hpp"
+
+namespace hiptensor::tuning::permutation
+{
+
+    void genInstances_F16_3_256_128_128_32_32(
+        std::vector<std::unique_ptr<
+            DeviceElementwise<ck::Tuple<F16>, ck::Tuple<F16>, UnaryScaleSquare, 3>>>& instances)
+    {
+        genInstances_256_128_128_32_32<F16, 3>(instances);
+    }
+
+}

tuning/permutation/permutation_instance_F16_3_256_128_128_8_8.cpp

@@ -0,0 +1,13 @@
+#include "permutation_instance.hpp"
+
+namespace hiptensor::tuning::permutation
+{
+
+    void genInstances_F16_3_256_128_128_8_8(
+        std::vector<std::unique_ptr<
+            DeviceElementwise<ck::Tuple<F16>, ck::Tuple<F16>, UnaryScaleSquare, 3>>>& instances)
+    {
+        genInstances_256_128_128_8_8<F16, 3>(instances);
+    }
+
+}

tuning/permutation/permutation_instance_F16_3_256_64_64_16_16.cpp

@@ -0,0 +1,13 @@
+#include "permutation_instance.hpp"
+
+namespace hiptensor::tuning::permutation
+{
+
+    void genInstances_F16_3_256_64_64_16_16(
+        std::vector<std::unique_ptr<
+            DeviceElementwise<ck::Tuple<F16>, ck::Tuple<F16>, UnaryScaleSquare, 3>>>& instances)
+    {
+        genInstances_256_64_64_16_16<F16, 3>(instances);
+    }
+
+}

tuning/permutation/permutation_instance_F16_3_256_64_64_4_4.cpp

@@ -0,0 +1,13 @@
+#include "permutation_instance.hpp"
+
+namespace hiptensor::tuning::permutation
+{
+
+    void genInstances_F16_3_256_64_64_4_4(
+        std::vector<std::unique_ptr<
+            DeviceElementwise<ck::Tuple<F16>, ck::Tuple<F16>, UnaryScaleSquare, 3>>>& instances)
+    {
+        genInstances_256_64_64_4_4<F16, 3>(instances);
+    }
+
+}

tuning/permutation/permutation_instance_F16_3_miscellaneous.cpp

@@ -0,0 +1,13 @@
+#include "permutation_instance.hpp"
+
+namespace hiptensor::tuning::permutation
+{
+
+    void genInstances_F16_3_miscellaneous(
+        std::vector<std::unique_ptr<
+            DeviceElementwise<ck::Tuple<F16>, ck::Tuple<F16>, UnaryScaleSquare, 3>>>& instances)
+    {
+        genInstances_miscellaneous<F16, 3>(instances);
+    }
+
+}

tuning/permutation/permutation_instance_F32_3_256_128_128_16_16.cpp

@@ -0,0 +1,13 @@
+#include "permutation_instance.hpp"
+
+namespace hiptensor::tuning::permutation
+{
+
+    void genInstances_F32_3_256_128_128_16_16(
+        std::vector<std::unique_ptr<
+            DeviceElementwise<ck::Tuple<F32>, ck::Tuple<F32>, UnaryScaleSquare, 3>>>& instances)
+    {
+        genInstances_256_128_128_16_16<F32, 3>(instances);
+    }
+
+}

tuning/permutation/permutation_instance_F32_3_256_128_128_32_32.cpp

@@ -0,0 +1,13 @@
+#include "permutation_instance.hpp"
+
+namespace hiptensor::tuning::permutation
+{
+
+    void genInstances_F32_3_256_128_128_32_32(
+        std::vector<std::unique_ptr<
+            DeviceElementwise<ck::Tuple<F32>, ck::Tuple<F32>, UnaryScaleSquare, 3>>>& instances)
+    {
+        genInstances_256_128_128_32_32<F32, 3>(instances);
+    }
+
+}

tuning/permutation/permutation_instance_F32_3_256_128_128_8_8.cpp

@@ -0,0 +1,13 @@
+#include "permutation_instance.hpp"
+
+namespace hiptensor::tuning::permutation
+{
+
+    void genInstances_F32_3_256_128_128_8_8(
+        std::vector<std::unique_ptr<
+            DeviceElementwise<ck::Tuple<F32>, ck::Tuple<F32>, UnaryScaleSquare, 3>>>& instances)
+    {
+        genInstances_256_128_128_8_8<F32, 3>(instances);
+    }
+
+}

tuning/permutation/permutation_instance_F32_3_256_64_64_16_16.cpp

@@ -0,0 +1,13 @@
+#include "permutation_instance.hpp"
+
+namespace hiptensor::tuning::permutation
+{
+
+    void genInstances_F32_3_256_64_64_16_16(
+        std::vector<std::unique_ptr<
+            DeviceElementwise<ck::Tuple<F32>, ck::Tuple<F32>, UnaryScaleSquare, 3>>>& instances)
+    {
+        genInstances_256_64_64_16_16<F32, 3>(instances);
+    }
+
+}

tuning/permutation/permutation_instance_F32_3_256_64_64_4_4.cpp

@@ -0,0 +1,13 @@
+#include "permutation_instance.hpp"
+
+namespace hiptensor::tuning::permutation
+{
+
+    void genInstances_F32_3_256_64_64_4_4(
+        std::vector<std::unique_ptr<
+            DeviceElementwise<ck::Tuple<F32>, ck::Tuple<F32>, UnaryScaleSquare, 3>>>& instances)
+    {
+        genInstances_256_64_64_4_4<F32, 3>(instances);
+    }
+
+}

tuning/permutation/permutation_instance_F32_3_miscellaneous.cpp

@@ -0,0 +1,13 @@
+#include "permutation_instance.hpp"
+
+namespace hiptensor::tuning::permutation
+{
+
+    void genInstances_F32_3_miscellaneous(
+        std::vector<std::unique_ptr<
+            DeviceElementwise<ck::Tuple<F32>, ck::Tuple<F32>, UnaryScaleSquare, 3>>>& instances)
+    {
+        genInstances_miscellaneous<F32, 3>(instances);
+    }
+
+}