add architecture

example/47_gemm_bias_softmax_gemm_permute/CMakeLists.txt

 add_example_executable(example_gemm_bias_softmax_gemm_permute gemm_bias_softmax_gemm_permute.cpp)
-add_example_executable(example_gemm_bias_softmax_gemm_permute_nolib gemm_bias_softmax_gemm_permute_nolib.cpp)

library/include/ck/library/tensor_operation_instance/add_device_operation_instance.hpp

@@ -29,7 +29,12 @@ void add_device_operation_instances(std::vector<std::unique_ptr<BaseOp>>& op_ins
     });
 }
 
-template <typename DeviceOp, typename Tag = void>
+enum struct ArchitectureEnum
+{
+    Xdl,
+    Dl
+};
+template <typename DeviceOp, ArchitectureEnum Arch = ArchitectureEnum::Xdl>
 struct DeviceOperationInstanceCreator;
 } // namespace instance
 } // namespace device

library/include/ck/library/tensor_operation_instance/gpu/batched_gemm_softmax_gemm_permute/batched_gemm_softmax_gemm_permute.hpp

@@ -34,7 +34,8 @@ template <index_t NumDimG,
           typename C0DEElementwiseOperation,
           typename B1ElementwiseOperation,
           typename C1DEElementwiseOperation,
-          MaskingSpecialization MaskingSpec>
+          MaskingSpecialization MaskingSpec,
+          ArchitectureEnum Arch = ArchitectureEnum::Xdl>
 void add_device_instances(
     std::vector<std::unique_ptr<DeviceBatchedGemmSoftmaxGemmPermute<NumDimG,
                                                                     NumDimM,
@@ -72,7 +73,7 @@ void add_device_instances(
                                                          C1DEElementwiseOperation,
                                                          MaskingSpec>;
     add_device_operation_instances(
-        instances, DeviceOperationInstanceCreator<DeviceOp>::create_device_instances());
+        instances, DeviceOperationInstanceCreator<DeviceOp, Arch>::create_device_instances());
 }
 } // namespace instance
 } // namespace device

library/include/ck/library/tensor_operation_instance/gpu/batched_gemm_softmax_gemm_permute/device_batched_gemm_multiple_d_softmax_gemm_permute_xdl_cshuffle_bf16_gmk_gnk_gno_gmo_instance.hpp

@@ -92,7 +92,8 @@ struct DeviceOperationInstanceCreator<DeviceBatchedGemmSoftmaxGemmPermute<NumDim
                                                                           C0DEElementwiseOperation,
                                                                           B1ElementwiseOperation,
                                                                           C1DEElementwiseOperation,
-                                                                          MaskingSpec>>
+                                                                          MaskingSpec>,
+                                      ArchitectureEnum::Xdl>
 {
     static auto create_device_instances()
     {

library/include/ck/library/tensor_operation_instance/gpu/batched_gemm_softmax_gemm_permute/device_batched_gemm_multiple_d_softmax_gemm_permute_xdl_cshuffle_fp16_gmk_gnk_gno_gmo_instance.hpp

@@ -92,7 +92,8 @@ struct DeviceOperationInstanceCreator<DeviceBatchedGemmSoftmaxGemmPermute<NumDim
                                                                           C0DEElementwiseOperation,
                                                                           B1ElementwiseOperation,
                                                                           C1DEElementwiseOperation,
-                                                                          MaskingSpec>>
+                                                                          MaskingSpec>,
+                                      ArchitectureEnum::Xdl>
 {
     static auto create_device_instances()
     {

src_example/01_gemm_bias_softmax_gemm_permute/gemm_bias_softmax_gemm_permute.cpp

@@ -135,7 +135,25 @@ int main()
 
     // get device op instances
     std::vector<std::unique_ptr<DeviceOp>> op_ptrs;
-    ck::tensor_operation::device::instance::add_device_instances(op_ptrs);
+    ck::tensor_operation::device::instance::add_device_instances<
+        2,
+        1,
+        1,
+        1,
+        1,
+        ADataType,
+        B0DataType,
+        B1DataType,
+        CDataType,
+        ck::Tuple<D00DataType, D01DataType>,
+        ck::Tuple<>,
+        AElementOp,
+        B0ElementOp,
+        Acc0ElementOp,
+        B1ElementOp,
+        CElementOp,
+        MaskingSpec,
+        ck::tensor_operation::device::instance::ArchitectureEnum::Xdl>(op_ptrs);
 
     std::cout << "found " << op_ptrs.size() << " instances" << std::endl;
 

src_example/CMakeLists.txt

+cmake_minimum_required(VERSION 3.15)
+project(ck_app)
+add_compile_options(-std=c++17)
+
+find_package(composable_kernel 1.0.0 COMPONENTS device_operations)
+find_package(hip REQUIRED PATHS /opt/rocm)
+message(STATUS "Build with HIP ${hip_VERSION}")
+
+# add all example subdir
+file(GLOB dir_list LIST_DIRECTORIES true *)
+FOREACH(subdir ${dir_list})
+    IF(IS_DIRECTORY "${subdir}" AND (NOT "${subdir}" MATCHES "build"))
+        add_subdirectory(${subdir})
+    ENDIF()
+ENDFOREACH()