Add instance for layout=TN

include/ck/ck.hpp

@@ -140,7 +140,7 @@
 // this will let make_default_loop_scheduler() return interwave scheduling flag by default
 #define CK_EXPERIMENTAL_DEFAULT_TO_INTER_WAVE_SCHEDULING 0
 // experimental feature: add instances using interwave scheduling
-#define CK_EXPERIMENTAL_INTER_WAVE_INSTANCES 1
+#define CK_EXPERIMENTAL_INTER_WAVE_INSTANCES 0
 // experimental feature: add instances using pipeline v2
 #define CK_EXPERIMENTAL_PIPELINE_V2_INSTANCES 1
 

include/ck/tensor_operation/gpu/grid/gridwise_gemm_pipeline_v2.hpp

@@ -79,6 +79,10 @@ struct GridwiseGemmPipeline_v2
 
             do
             {
+                #ifdef USE_IGLP_OPT
+                __builtin_amdgcn_iglp_opt(1);
+                #endif
+
                 block_sync_lds();
 
                 // GEMM i

include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v2r3.hpp

@@ -27,6 +27,9 @@ template <typename GridwiseGemm,
           typename CElementwiseOperation,
           typename Block2CTileMap,
           bool HasMainKBlockLoop>
+#ifdef USE_WAVES_PER_EU
+__attribute__((amdgpu_waves_per_eu(1, 1)))
+#endif
 __global__ void
 #if CK_USE_LAUNCH_BOUNDS
     __launch_bounds__(CK_MAX_THREAD_PER_BLOCK, CK_MIN_BLOCK_PER_CU)

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

@@ -257,123 +257,122 @@ struct DeviceOperationInstanceFactory<
     {
         std::vector<std::unique_ptr<DeviceOp>> op_ptrs;
 
-        if constexpr(is_same_v<ADataType, float> && is_same_v<BDataType, float> &&
-                     is_same_v<CDataType, float>)
-        {
-            if constexpr(is_same_v<ALayout, Row> && is_same_v<BLayout, Row> &&
-                         is_same_v<CLayout, Row>)
-            {
-                add_device_gemm_xdl_f32_f32_f32_mk_kn_mn_instances(op_ptrs);
-                add_device_gemm_dl_f32_f32_f32_mk_kn_mn_instances(op_ptrs);
-                add_device_gemm_xdl_c_shuffle_f32_f32_f32_mk_kn_mn_instances(op_ptrs);
-            }
-            else if constexpr(is_same_v<ALayout, Row> && is_same_v<BLayout, Col> &&
-                              is_same_v<CLayout, Row>)
-            {
-                add_device_gemm_xdl_f32_f32_f32_mk_nk_mn_instances(op_ptrs);
-                add_device_gemm_dl_f32_f32_f32_mk_nk_mn_instances(op_ptrs);
-                add_device_gemm_xdl_c_shuffle_f32_f32_f32_mk_nk_mn_instances(op_ptrs);
-            }
-            else if constexpr(is_same_v<ALayout, Col> && is_same_v<BLayout, Row> &&
-                              is_same_v<CLayout, Row>)
-            {
-                add_device_gemm_xdl_f32_f32_f32_km_kn_mn_instances(op_ptrs);
-                add_device_gemm_dl_f32_f32_f32_km_kn_mn_instances(op_ptrs);
-                add_device_gemm_xdl_c_shuffle_f32_f32_f32_km_kn_mn_instances(op_ptrs);
-            }
-            else if constexpr(is_same_v<ALayout, Col> && is_same_v<BLayout, Col> &&
-                              is_same_v<CLayout, Row>)
-            {
-                add_device_gemm_xdl_f32_f32_f32_km_nk_mn_instances(op_ptrs);
-                add_device_gemm_dl_f32_f32_f32_km_nk_mn_instances(op_ptrs);
-                add_device_gemm_xdl_c_shuffle_f32_f32_f32_km_nk_mn_instances(op_ptrs);
-            }
-        }
-        else if constexpr(is_same_v<ADataType, half_t> && is_same_v<BDataType, half_t> &&
+        // if constexpr(is_same_v<ADataType, float> && is_same_v<BDataType, float> &&
+        //              is_same_v<CDataType, float>)
+        // {
+        //     if constexpr(is_same_v<ALayout, Row> && is_same_v<BLayout, Row> &&
+        //                  is_same_v<CLayout, Row>)
+        //     {
+        //         add_device_gemm_xdl_f32_f32_f32_mk_kn_mn_instances(op_ptrs);
+        //         add_device_gemm_dl_f32_f32_f32_mk_kn_mn_instances(op_ptrs);
+        //         add_device_gemm_xdl_c_shuffle_f32_f32_f32_mk_kn_mn_instances(op_ptrs);
+        //     }
+        //     else if constexpr(is_same_v<ALayout, Row> && is_same_v<BLayout, Col> &&
+        //                       is_same_v<CLayout, Row>)
+        //     {
+        //         add_device_gemm_xdl_f32_f32_f32_mk_nk_mn_instances(op_ptrs);
+        //         add_device_gemm_dl_f32_f32_f32_mk_nk_mn_instances(op_ptrs);
+        //         add_device_gemm_xdl_c_shuffle_f32_f32_f32_mk_nk_mn_instances(op_ptrs);
+        //     }
+        //     else if constexpr(is_same_v<ALayout, Col> && is_same_v<BLayout, Row> &&
+        //                       is_same_v<CLayout, Row>)
+        //     {
+        //         add_device_gemm_xdl_f32_f32_f32_km_kn_mn_instances(op_ptrs);
+        //         add_device_gemm_dl_f32_f32_f32_km_kn_mn_instances(op_ptrs);
+        //         add_device_gemm_xdl_c_shuffle_f32_f32_f32_km_kn_mn_instances(op_ptrs);
+        //     }
+        //     else if constexpr(is_same_v<ALayout, Col> && is_same_v<BLayout, Col> &&
+        //                       is_same_v<CLayout, Row>)
+        //     {
+        //         add_device_gemm_xdl_f32_f32_f32_km_nk_mn_instances(op_ptrs);
+        //         add_device_gemm_dl_f32_f32_f32_km_nk_mn_instances(op_ptrs);
+        //         add_device_gemm_xdl_c_shuffle_f32_f32_f32_km_nk_mn_instances(op_ptrs);
+        //     }
+        // }
+        if constexpr(is_same_v<ADataType, half_t> && is_same_v<BDataType, half_t> &&
                           is_same_v<CDataType, half_t>)
         {
             if constexpr(is_same_v<ALayout, Row> && is_same_v<BLayout, Row> &&
                          is_same_v<CLayout, Row>)
             {
-                add_device_gemm_xdl_f16_f16_f16_mk_kn_mn_instances(op_ptrs);
-                add_device_gemm_dl_f16_f16_f16_mk_kn_mn_instances(op_ptrs);
-                add_device_gemm_xdl_c_shuffle_f16_f16_f16_mk_kn_mn_instances(op_ptrs);
+                // add_device_gemm_xdl_f16_f16_f16_mk_kn_mn_instances(op_ptrs);
+                // add_device_gemm_dl_f16_f16_f16_mk_kn_mn_instances(op_ptrs);
+                // add_device_gemm_xdl_c_shuffle_f16_f16_f16_mk_kn_mn_instances(op_ptrs);
             }
             else if constexpr(is_same_v<ALayout, Row> && is_same_v<BLayout, Col> &&
                               is_same_v<CLayout, Row>)
             {
                 add_device_gemm_xdl_f16_f16_f16_mk_nk_mn_instances(op_ptrs);
-                add_device_gemm_dl_f16_f16_f16_mk_nk_mn_instances(op_ptrs);
-                add_device_gemm_xdl_c_shuffle_f16_f16_f16_mk_nk_mn_instances(op_ptrs);
-                add_device_gemm_xdl_c_shuffle_2_stage_f16_f16_f16_mk_nk_mn_instances(op_ptrs);
-            }
-            else if constexpr(is_same_v<ALayout, Col> && is_same_v<BLayout, Row> &&
-                              is_same_v<CLayout, Row>)
-            {
-                add_device_gemm_xdl_f16_f16_f16_km_kn_mn_instances(op_ptrs);
-                add_device_gemm_dl_f16_f16_f16_km_kn_mn_instances(op_ptrs);
-                add_device_gemm_xdl_c_shuffle_f16_f16_f16_km_kn_mn_instances(op_ptrs);
-            }
-            else if constexpr(is_same_v<ALayout, Col> && is_same_v<BLayout, Col> &&
-                              is_same_v<CLayout, Row>)
-            {
-                add_device_gemm_xdl_f16_f16_f16_km_nk_mn_instances(op_ptrs);
-                add_device_gemm_dl_f16_f16_f16_km_nk_mn_instances(op_ptrs);
-                add_device_gemm_xdl_c_shuffle_f16_f16_f16_km_nk_mn_instances(op_ptrs);
-            }
-        }
-        else if constexpr(is_same_v<ADataType, ck::bhalf_t> && is_same_v<BDataType, ck::bhalf_t> &&
-                          is_same_v<CDataType, ck::bhalf_t>)
-        {
-            if constexpr(is_same_v<ALayout, Row> && is_same_v<BLayout, Row> &&
-                         is_same_v<CLayout, Row>)
-            {
-                add_device_gemm_xdl_c_shuffle_bf16_bf16_bf16_mk_kn_mn_instances(op_ptrs);
-            }
-            else if constexpr(is_same_v<ALayout, Row> && is_same_v<BLayout, Col> &&
-                              is_same_v<CLayout, Row>)
-            {
-                add_device_gemm_xdl_c_shuffle_bf16_bf16_bf16_mk_nk_mn_instances(op_ptrs);
-            }
-            else if constexpr(is_same_v<ALayout, Col> && is_same_v<BLayout, Row> &&
-                              is_same_v<CLayout, Row>)
-            {
-                add_device_gemm_xdl_c_shuffle_bf16_bf16_bf16_km_kn_mn_instances(op_ptrs);
-            }
-            else if constexpr(is_same_v<ALayout, Col> && is_same_v<BLayout, Col> &&
-                              is_same_v<CLayout, Row>)
-            {
-                add_device_gemm_xdl_c_shuffle_bf16_bf16_bf16_km_nk_mn_instances(op_ptrs);
-            }
-        }
-        else if constexpr(is_same_v<ADataType, int8_t> && is_same_v<BDataType, int8_t> &&
-                          is_same_v<CDataType, int8_t>)
-        {
-            if constexpr(is_same_v<ALayout, Row> && is_same_v<BLayout, Row> &&
-                         is_same_v<CLayout, Row>)
-            {
-                add_device_gemm_xdl_c_shuffle_i8_i8_i8_mk_kn_mn_instances(op_ptrs);
-                add_device_gemm_dl_i8_i8_i8_mk_kn_mn_instances(op_ptrs);
-            }
-            else if constexpr(is_same_v<ALayout, Row> && is_same_v<BLayout, Col> &&
-                              is_same_v<CLayout, Row>)
-            {
-                add_device_gemm_xdl_c_shuffle_i8_i8_i8_mk_nk_mn_instances(op_ptrs);
-                add_device_gemm_dl_i8_i8_i8_mk_nk_mn_instances(op_ptrs);
+                // add_device_gemm_xdl_c_shuffle_f16_f16_f16_mk_nk_mn_instances(op_ptrs);
+                // add_device_gemm_xdl_c_shuffle_2_stage_f16_f16_f16_mk_nk_mn_instances(op_ptrs);
             }
             else if constexpr(is_same_v<ALayout, Col> && is_same_v<BLayout, Row> &&
                               is_same_v<CLayout, Row>)
             {
-                add_device_gemm_xdl_c_shuffle_i8_i8_i8_km_kn_mn_instances(op_ptrs);
-                add_device_gemm_dl_i8_i8_i8_km_kn_mn_instances(op_ptrs);
+                // add_device_gemm_xdl_f16_f16_f16_km_kn_mn_instances(op_ptrs);
+                // add_device_gemm_dl_f16_f16_f16_km_kn_mn_instances(op_ptrs);
+                // add_device_gemm_xdl_c_shuffle_f16_f16_f16_km_kn_mn_instances(op_ptrs);
             }
             else if constexpr(is_same_v<ALayout, Col> && is_same_v<BLayout, Col> &&
                               is_same_v<CLayout, Row>)
             {
-                add_device_gemm_xdl_c_shuffle_i8_i8_i8_km_nk_mn_instances(op_ptrs);
-                add_device_gemm_dl_i8_i8_i8_km_nk_mn_instances(op_ptrs);
+                // add_device_gemm_xdl_f16_f16_f16_km_nk_mn_instances(op_ptrs);
+                // add_device_gemm_dl_f16_f16_f16_km_nk_mn_instances(op_ptrs);
+                // add_device_gemm_xdl_c_shuffle_f16_f16_f16_km_nk_mn_instances(op_ptrs);
             }
         }
+        // else if constexpr(is_same_v<ADataType, ck::bhalf_t> && is_same_v<BDataType, ck::bhalf_t> &&
+        //                   is_same_v<CDataType, ck::bhalf_t>)
+        // {
+        //     if constexpr(is_same_v<ALayout, Row> && is_same_v<BLayout, Row> &&
+        //                  is_same_v<CLayout, Row>)
+        //     {
+        //         add_device_gemm_xdl_c_shuffle_bf16_bf16_bf16_mk_kn_mn_instances(op_ptrs);
+        //     }
+        //     else if constexpr(is_same_v<ALayout, Row> && is_same_v<BLayout, Col> &&
+        //                       is_same_v<CLayout, Row>)
+        //     {
+        //         add_device_gemm_xdl_c_shuffle_bf16_bf16_bf16_mk_nk_mn_instances(op_ptrs);
+        //     }
+        //     else if constexpr(is_same_v<ALayout, Col> && is_same_v<BLayout, Row> &&
+        //                       is_same_v<CLayout, Row>)
+        //     {
+        //         add_device_gemm_xdl_c_shuffle_bf16_bf16_bf16_km_kn_mn_instances(op_ptrs);
+        //     }
+        //     else if constexpr(is_same_v<ALayout, Col> && is_same_v<BLayout, Col> &&
+        //                       is_same_v<CLayout, Row>)
+        //     {
+        //         add_device_gemm_xdl_c_shuffle_bf16_bf16_bf16_km_nk_mn_instances(op_ptrs);
+        //     }
+        // }
+        // else if constexpr(is_same_v<ADataType, int8_t> && is_same_v<BDataType, int8_t> &&
+        //                   is_same_v<CDataType, int8_t>)
+        // {
+        //     if constexpr(is_same_v<ALayout, Row> && is_same_v<BLayout, Row> &&
+        //                  is_same_v<CLayout, Row>)
+        //     {
+        //         add_device_gemm_xdl_c_shuffle_i8_i8_i8_mk_kn_mn_instances(op_ptrs);
+        //         add_device_gemm_dl_i8_i8_i8_mk_kn_mn_instances(op_ptrs);
+        //     }
+        //     else if constexpr(is_same_v<ALayout, Row> && is_same_v<BLayout, Col> &&
+        //                       is_same_v<CLayout, Row>)
+        //     {
+        //         add_device_gemm_xdl_c_shuffle_i8_i8_i8_mk_nk_mn_instances(op_ptrs);
+        //         add_device_gemm_dl_i8_i8_i8_mk_nk_mn_instances(op_ptrs);
+        //     }
+        //     else if constexpr(is_same_v<ALayout, Col> && is_same_v<BLayout, Row> &&
+        //                       is_same_v<CLayout, Row>)
+        //     {
+        //         add_device_gemm_xdl_c_shuffle_i8_i8_i8_km_kn_mn_instances(op_ptrs);
+        //         add_device_gemm_dl_i8_i8_i8_km_kn_mn_instances(op_ptrs);
+        //     }
+        //     else if constexpr(is_same_v<ALayout, Col> && is_same_v<BLayout, Col> &&
+        //                       is_same_v<CLayout, Row>)
+        //     {
+        //         add_device_gemm_xdl_c_shuffle_i8_i8_i8_km_nk_mn_instances(op_ptrs);
+        //         add_device_gemm_dl_i8_i8_i8_km_nk_mn_instances(op_ptrs);
+        //     }
+        // }
 
         return op_ptrs;
     }

library/src/tensor_operation_instance/gpu/gemm/CMakeLists.txt

 add_instance_library(device_gemm_instance
-   device_gemm_xdl_f64_f64_f64_mk_kn_mn_instance.cpp
-   device_gemm_xdl_f64_f64_f64_mk_nk_mn_instance.cpp
-   device_gemm_xdl_f64_f64_f64_km_kn_mn_instance.cpp
-   device_gemm_xdl_f64_f64_f64_km_nk_mn_instance.cpp
-   device_gemm_xdl_f32_f32_f32_mk_kn_mn_instance.cpp
-   device_gemm_xdl_f32_f32_f32_mk_nk_mn_instance.cpp
-   device_gemm_xdl_f32_f32_f32_km_kn_mn_instance.cpp
-   device_gemm_xdl_f32_f32_f32_km_nk_mn_instance.cpp
-   device_gemm_xdl_f16_f16_f16_mk_kn_mn_instance.cpp
+   # device_gemm_xdl_f64_f64_f64_mk_kn_mn_instance.cpp
+   # device_gemm_xdl_f64_f64_f64_mk_nk_mn_instance.cpp
+   # device_gemm_xdl_f64_f64_f64_km_kn_mn_instance.cpp
+   # device_gemm_xdl_f64_f64_f64_km_nk_mn_instance.cpp
+   # device_gemm_xdl_f32_f32_f32_mk_kn_mn_instance.cpp
+   # device_gemm_xdl_f32_f32_f32_mk_nk_mn_instance.cpp
+   # device_gemm_xdl_f32_f32_f32_km_kn_mn_instance.cpp
+   # device_gemm_xdl_f32_f32_f32_km_nk_mn_instance.cpp
+   # device_gemm_xdl_f16_f16_f16_mk_kn_mn_instance.cpp
    device_gemm_xdl_f16_f16_f16_mk_nk_mn_instance.cpp
-   device_gemm_xdl_f16_f16_f16_km_kn_mn_instance.cpp
-   device_gemm_xdl_f16_f16_f16_km_nk_mn_instance.cpp
-   device_gemm_xdl_c_shuffle_i8_i8_i8_mk_kn_mn_instance.cpp
-   device_gemm_xdl_c_shuffle_i8_i8_i8_mk_nk_mn_instance.cpp
-   device_gemm_xdl_c_shuffle_i8_i8_i8_km_kn_mn_instance.cpp
-   device_gemm_xdl_c_shuffle_i8_i8_i8_km_nk_mn_instance.cpp
-   device_gemm_xdl_c_shuffle_bf16_bf16_bf16_mk_kn_mn_instance.cpp
-   device_gemm_xdl_c_shuffle_bf16_bf16_bf16_mk_nk_mn_instance.cpp
-   device_gemm_xdl_c_shuffle_bf16_bf16_bf16_km_kn_mn_instance.cpp
-   device_gemm_xdl_c_shuffle_bf16_bf16_bf16_km_nk_mn_instance.cpp
-   device_gemm_xdl_c_shuffle_f16_f16_f16_mk_kn_mn_instance.cpp
-   device_gemm_xdl_c_shuffle_f16_f16_f16_mk_nk_mn_instance.cpp
-   device_gemm_xdl_c_shuffle_f16_f16_f16_km_kn_mn_instance.cpp
-   device_gemm_xdl_c_shuffle_f16_f16_f16_km_nk_mn_instance.cpp
-   device_gemm_xdl_c_shuffle_f32_f32_f32_mk_kn_mn_instance.cpp
-   device_gemm_xdl_c_shuffle_f32_f32_f32_mk_nk_mn_instance.cpp
-   device_gemm_xdl_c_shuffle_f32_f32_f32_km_kn_mn_instance.cpp
-   device_gemm_xdl_c_shuffle_f32_f32_f32_km_nk_mn_instance.cpp
-   device_gemm_xdl_c_shuffle_2_stage_f16_f16_f16_mk_nk_mn_instance.cpp
-   device_gemm_dl_f32_f32_f32_mk_kn_mn_instance.cpp
-   device_gemm_dl_f32_f32_f32_mk_nk_mn_instance.cpp
-   device_gemm_dl_f32_f32_f32_km_kn_mn_instance.cpp
-   device_gemm_dl_f32_f32_f32_km_nk_mn_instance.cpp
-   device_gemm_dl_f16_f16_f16_mk_kn_mn_instance.cpp
-   device_gemm_dl_f16_f16_f16_mk_nk_mn_instance.cpp
-   device_gemm_dl_f16_f16_f16_km_kn_mn_instance.cpp
-   device_gemm_dl_f16_f16_f16_km_nk_mn_instance.cpp
-   device_gemm_dl_i8_i8_i8_mk_kn_mn_instance.cpp
-   device_gemm_dl_i8_i8_i8_mk_nk_mn_instance.cpp
-   device_gemm_dl_i8_i8_i8_km_kn_mn_instance.cpp
-   device_gemm_dl_i8_i8_i8_km_nk_mn_instance.cpp
+   # device_gemm_xdl_f16_f16_f16_km_kn_mn_instance.cpp
+   # device_gemm_xdl_f16_f16_f16_km_nk_mn_instance.cpp
+   # device_gemm_xdl_c_shuffle_i8_i8_i8_mk_kn_mn_instance.cpp
+   # device_gemm_xdl_c_shuffle_i8_i8_i8_mk_nk_mn_instance.cpp
+   # device_gemm_xdl_c_shuffle_i8_i8_i8_km_kn_mn_instance.cpp
+   # device_gemm_xdl_c_shuffle_i8_i8_i8_km_nk_mn_instance.cpp
+   # device_gemm_xdl_c_shuffle_bf16_bf16_bf16_mk_kn_mn_instance.cpp
+   # device_gemm_xdl_c_shuffle_bf16_bf16_bf16_mk_nk_mn_instance.cpp
+   # device_gemm_xdl_c_shuffle_bf16_bf16_bf16_km_kn_mn_instance.cpp
+   # device_gemm_xdl_c_shuffle_bf16_bf16_bf16_km_nk_mn_instance.cpp
+   # device_gemm_xdl_c_shuffle_f16_f16_f16_mk_kn_mn_instance.cpp
+   ## device_gemm_xdl_c_shuffle_f16_f16_f16_mk_nk_mn_instance.cpp
+   # device_gemm_xdl_c_shuffle_f16_f16_f16_km_kn_mn_instance.cpp
+   # device_gemm_xdl_c_shuffle_f16_f16_f16_km_nk_mn_instance.cpp
+   # device_gemm_xdl_c_shuffle_f32_f32_f32_mk_kn_mn_instance.cpp
+   # device_gemm_xdl_c_shuffle_f32_f32_f32_mk_nk_mn_instance.cpp
+   # device_gemm_xdl_c_shuffle_f32_f32_f32_km_kn_mn_instance.cpp
+   # device_gemm_xdl_c_shuffle_f32_f32_f32_km_nk_mn_instance.cpp
+   ## device_gemm_xdl_c_shuffle_2_stage_f16_f16_f16_mk_nk_mn_instance.cpp
+   # device_gemm_dl_f32_f32_f32_mk_kn_mn_instance.cpp
+   # device_gemm_dl_f32_f32_f32_mk_nk_mn_instance.cpp
+   # device_gemm_dl_f32_f32_f32_km_kn_mn_instance.cpp
+   # device_gemm_dl_f32_f32_f32_km_nk_mn_instance.cpp
+   # device_gemm_dl_f16_f16_f16_mk_kn_mn_instance.cpp
+   ## device_gemm_dl_f16_f16_f16_mk_nk_mn_instance.cpp
+   # device_gemm_dl_f16_f16_f16_km_kn_mn_instance.cpp
+   # device_gemm_dl_f16_f16_f16_km_nk_mn_instance.cpp
+   # device_gemm_dl_i8_i8_i8_mk_kn_mn_instance.cpp
+   # device_gemm_dl_i8_i8_i8_mk_nk_mn_instance.cpp
+   # device_gemm_dl_i8_i8_i8_km_kn_mn_instance.cpp
+   # device_gemm_dl_i8_i8_i8_km_nk_mn_instance.cpp
 )
+
+set(ENABLE_IGLP_OPT ON)
+if(ENABLE_IGLP_OPT)
+set_source_files_properties(device_gemm_xdl_f16_f16_f16_mk_nk_mn_instance.cpp PROPERTIES
+  COMPILE_OPTIONS ";--save-temps;-Wno-gnu-line-marker;-mllvm;-amdgpu-enable-max-ilp-scheduling-strategy;"
+  COMPILE_DEFINITIONS ";USE_IGLP_OPT;USE_WAVES_PER_EU;")
+endif(ENABLE_IGLP_OPT)

profiler/src/CMakeLists.txt

@@ -2,30 +2,30 @@
 set(PROFILER_SOURCES
     profiler.cpp
     profile_gemm.cpp
-    profile_gemm_splitk.cpp
-    profile_gemm_bilinear.cpp
-    profile_gemm_bias_add_reduce.cpp
-    profile_gemm_add_add_fastgelu.cpp
-    profile_gemm_add_fastgelu.cpp
-    profile_gemm_fastgelu.cpp
-    profile_gemm_reduce.cpp
-    profile_batched_gemm.cpp
-    profile_batched_gemm_gemm.cpp
-    profile_batched_gemm_add_relu_gemm_add.cpp
-    profile_batched_gemm_reduce.cpp
-    profile_grouped_gemm.cpp
-    profile_conv_fwd.cpp
-    profile_conv_fwd_bias_relu.cpp
-    profile_conv_fwd_bias_relu_add.cpp
-    profile_conv_bwd_data.cpp
-    profile_grouped_conv_fwd.cpp
-    profile_grouped_conv_bwd_weight.cpp
-    profile_reduce.cpp
-    profile_groupnorm.cpp
-    profile_layernorm.cpp
-    profile_softmax.cpp
-    profile_batchnorm_fwd.cpp
-    profile_batchnorm_bwd.cpp
+    # profile_gemm_splitk.cpp
+    # profile_gemm_bilinear.cpp
+    # profile_gemm_bias_add_reduce.cpp
+    # profile_gemm_add_add_fastgelu.cpp
+    # profile_gemm_add_fastgelu.cpp
+    # profile_gemm_fastgelu.cpp
+    # profile_gemm_reduce.cpp
+    # profile_batched_gemm.cpp
+    # profile_batched_gemm_gemm.cpp
+    # profile_batched_gemm_add_relu_gemm_add.cpp
+    # profile_batched_gemm_reduce.cpp
+    # profile_grouped_gemm.cpp
+    # profile_conv_fwd.cpp
+    # profile_conv_fwd_bias_relu.cpp
+    # profile_conv_fwd_bias_relu_add.cpp
+    # profile_conv_bwd_data.cpp
+    # profile_grouped_conv_fwd.cpp
+    # profile_grouped_conv_bwd_weight.cpp
+    # profile_reduce.cpp
+    # profile_groupnorm.cpp
+    # profile_layernorm.cpp
+    # profile_softmax.cpp
+    # profile_batchnorm_fwd.cpp
+    # profile_batchnorm_bwd.cpp
 )
 
 set(PROFILER_EXECUTABLE ckProfiler)
@@ -35,33 +35,33 @@ target_compile_options(${PROFILER_EXECUTABLE} PRIVATE -Wno-global-constructors)
 
 target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE utility)
 target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_instance)
-target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_splitk_instance)
-target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_bilinear_instance)
-target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_add_add_fastgelu_instance)
-target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_add_fastgelu_instance)
-target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_fastgelu_instance)
-target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_reduce_instance)
-target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_bias_add_reduce_instance)
-target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_batched_gemm_instance)
-target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_batched_gemm_gemm_instance)
-target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_batched_gemm_add_relu_gemm_add_instance)
-target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_batched_gemm_reduce_instance)
-target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_grouped_gemm_instance)
-target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_conv2d_fwd_instance)
-target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_grouped_conv1d_fwd_instance)
-target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_grouped_conv2d_fwd_instance)
-target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_grouped_conv3d_fwd_instance)
-target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_conv1d_bwd_data_instance)
-target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_conv2d_bwd_data_instance)
-target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_conv3d_bwd_data_instance)
-target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_grouped_conv1d_bwd_weight_instance)
-target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_grouped_conv2d_bwd_weight_instance)
-target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_grouped_conv3d_bwd_weight_instance)
-target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_conv2d_fwd_bias_relu_instance)
-target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_conv2d_fwd_bias_relu_add_instance)
-target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_normalization_instance)
-target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_softmax_instance)
-target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_reduce_instance)
-target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_batchnorm_instance)
+# target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_splitk_instance)
+# target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_bilinear_instance)
+# target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_add_add_fastgelu_instance)
+# target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_add_fastgelu_instance)
+# target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_fastgelu_instance)
+# target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_reduce_instance)
+# target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_bias_add_reduce_instance)
+# target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_batched_gemm_instance)
+# target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_batched_gemm_gemm_instance)
+# target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_batched_gemm_add_relu_gemm_add_instance)
+# target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_batched_gemm_reduce_instance)
+# target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_grouped_gemm_instance)
+# target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_conv2d_fwd_instance)
+# target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_grouped_conv1d_fwd_instance)
+# target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_grouped_conv2d_fwd_instance)
+# target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_grouped_conv3d_fwd_instance)
+# target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_conv1d_bwd_data_instance)
+# target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_conv2d_bwd_data_instance)
+# target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_conv3d_bwd_data_instance)
+# target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_grouped_conv1d_bwd_weight_instance)
+# target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_grouped_conv2d_bwd_weight_instance)
+# target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_grouped_conv3d_bwd_weight_instance)
+# target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_conv2d_fwd_bias_relu_instance)
+# target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_conv2d_fwd_bias_relu_add_instance)
+# target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_normalization_instance)
+# target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_softmax_instance)
+# target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_reduce_instance)
+# target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_batchnorm_instance)
 
 rocm_install(TARGETS ${PROFILER_EXECUTABLE} COMPONENT profiler)