Fix Navi compilation.

include/ck/tensor_operation/gpu/device/impl/device_grouped_gemm_multiple_d_splitk_xdl_cshuffle_tile_loop.hpp

@@ -215,6 +215,9 @@ __global__ void
     ignore = p_workspace;
     ignore = tile_count;
     ignore = k_batch;
+    ignore = a_element_op;
+    ignore = b_element_op;
+    ignore = cde_element_op;
 #endif // end of if (defined(__gfx908__) || defined(__gfx90a__))
 }
 

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

@@ -16,6 +16,7 @@ namespace tensor_operation {
 namespace device {
 namespace instance {
 
+#if defined(CK_ENABLE_FP16)
 // MultiD version
 void add_device_grouped_gemm_multi_d_splitk_cshuffle_f16_f16_f16_mk_nk_mn_irregular_instances(
     std::vector<std::unique_ptr<DeviceGroupedGemm<Row,
@@ -42,6 +43,7 @@ void add_device_grouped_gemm_multi_d_splitk_cshuffle_f16_f16_f16_mk_kn_mn_irregu
                                                   PassThrough,
                                                   PassThrough,
                                                   PassThrough>>>& instances);
+#endif
 
 template <typename ALayout,
           typename BLayout,
@@ -87,14 +89,18 @@ struct DeviceOperationInstanceFactory<
             if constexpr(is_same_v<ALayout, Row> && is_same_v<BLayout, Row> &&
                          is_same_v<ELayout, Row>)
             {
+#if defined(CK_ENABLE_FP16)
                 add_device_grouped_gemm_multi_d_splitk_cshuffle_f16_f16_f16_mk_kn_mn_irregular_instances(
                     op_ptrs);
+#endif
             }
             else if constexpr(is_same_v<ALayout, Row> && is_same_v<BLayout, Col> &&
                               is_same_v<ELayout, Row>)
             {
+#if defined(CK_ENABLE_FP16)
                 add_device_grouped_gemm_multi_d_splitk_cshuffle_f16_f16_f16_mk_nk_mn_irregular_instances(
                     op_ptrs);
+#endif
             }
         }
         return op_ptrs;