Add tuning tool for permutation rank3

ROCm · Nov 20, 2024 · a3e2b0a · a3e2b0a
1 parent c97246a
commit a3e2b0a
Show file tree

Hide file tree

Showing 20 changed files with 1,108 additions and 152 deletions.
diff --git a/library/src/permutation/permutation_solution_impl.hpp b/library/src/permutation/permutation_solution_impl.hpp
@@ -181,40 +181,15 @@ namespace hiptensor
             ck::tensor_operation::element_wise::UnaryCombinedOp<Aop, Scale, Bop>,
             NumDim>;
 
-        using DeviceElementwisePermuteInstance
-            = ck::tensor_operation::device::DeviceElementwiseImpl<
-                InDataTypeTuple, // InDataTypeTuple
-                OutDataTypeTuple, // OutDataTypeTuple
-                // PassThrough,          // Elementwise
-                ck::tensor_operation::element_wise::UnaryCombinedOp<Aop, Scale, Bop>,
-                NumDim, // NumDim
-                256, // BlockSize
-                128, // M0PerBlock
-                128, // M1PerBlock
-                16, // M0PerThread
-                16, // M1PerThread
-                ck::Sequence<0, 1>, // ThreadClusterArrangeOrder
-                ck::Sequence<16>, // InScalarPerVectorSeq
-                ck::Sequence<16>>; // OutScalarPerVectorSeq
-
-        std::vector<std::unique_ptr<PermutationSolution>> result;
-        result.push_back(std::make_unique<PermutationSolutionImpl<PermutationOp>>(
-            std::make_unique<DeviceElementwisePermuteInstance>(
-                DeviceElementwisePermuteInstance{})));
-
-        /*
         using Factory
             = ck::tensor_operation::device::instance::DeviceOperationInstanceFactory<PermutationOp>;
 
         std::vector<std::unique_ptr<PermutationSolution>> result;
-		std::cout << "\n";
         for(auto& opPtr : Factory::GetInstances())
         {
             result.push_back(
                 std::make_unique<PermutationSolutionImpl<PermutationOp>>(std::move(opPtr)));
-			std::cout << result.back()->kernelName() << "\n";
         }
-*/
         return result;
     }
 

diff --git a/samples/CMakeLists.txt b/samples/CMakeLists.txt
@@ -44,7 +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/src/include/
                                 ${PROJECT_SOURCE_DIR}/library/include)
 
         # Build this sample under custom target

diff --git a/tuning/permutation/CMakeLists.txt b/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)