Compiler DL & XDL examples only when appropriate arch is available.

example/15_grouped_gemm/CMakeLists.txt

-add_custom_target(example_grouped_gemm_xdl)
-add_example_executable(example_grouped_gemm_xdl_fp32 grouped_gemm_xdl_fp32.cpp)
-add_example_dependencies(example_grouped_gemm_xdl example_grouped_gemm_xdl_fp32)
-
-add_example_executable(example_grouped_gemm_xdl_fp16 grouped_gemm_xdl_fp16.cpp)
-add_example_dependencies(example_grouped_gemm_xdl example_grouped_gemm_xdl_fp16)
+add_custom_target(example_grouped_gemm)
 
+if (DL_KERNELS)
 add_example_executable(example_grouped_gemm_multiple_d_dl_fp16 grouped_gemm_multiple_d_dl_fp16.cpp)
-add_example_dependencies(example_grouped_gemm_xdl example_grouped_gemm_multiple_d_dl_fp16)
+add_example_dependencies(example_grouped_gemm example_grouped_gemm_multiple_d_dl_fp16)
+endif()
+
+list(APPEND gpu_list gfx90a gfx940 gfx941 gfx942)
+set(target 0)
+foreach(gpu IN LISTS GPU_TARGETS)
+    if(gpu IN_LIST gpu_list AND target EQUAL 0)
+        add_example_executable(example_grouped_gemm_xdl_fp32 grouped_gemm_xdl_fp32.cpp)
+        add_example_dependencies(example_grouped_gemm example_grouped_gemm_xdl_fp32)
 
-add_example_executable(example_grouped_gemm_xdl_splitk_fp16 grouped_gemm_xdl_splitk_fp16.cpp)
-add_example_dependencies(example_grouped_gemm_xdl example_grouped_gemm_xdl_splitk_fp16)
+        add_example_executable(example_grouped_gemm_xdl_fp16 grouped_gemm_xdl_fp16.cpp)
+        add_example_dependencies(example_grouped_gemm example_grouped_gemm_xdl_fp16)
 
-add_example_executable(example_grouped_gemm_xdl_fixed_nk_fp16 grouped_gemm_xdl_fixed_nk_fp16.cpp)
-add_example_dependencies(example_grouped_gemm_xdl example_grouped_gemm_xdl_fixed_nk_fp16)
+        add_example_executable(example_grouped_gemm_xdl_splitk_fp16 grouped_gemm_xdl_splitk_fp16.cpp)
+        add_example_dependencies(example_grouped_gemm example_grouped_gemm_xdl_splitk_fp16)
 
-add_example_executable(example_grouped_gemm_xdl_fixed_nk_bias_fp16 grouped_gemm_xdl_fixed_nk_bias_fp16.cpp)
-add_example_dependencies(example_grouped_gemm_xdl example_grouped_gemm_xdl_fixed_nk_bias_fp16)
+        add_example_executable(example_grouped_gemm_xdl_fixed_nk_fp16 grouped_gemm_xdl_fixed_nk_fp16.cpp)
+        add_example_dependencies(example_grouped_gemm example_grouped_gemm_xdl_fixed_nk_fp16)
 
-add_example_executable(example_grouped_gemm_xdl_bf16 grouped_gemm_xdl_bf16.cpp)
-add_example_dependencies(example_grouped_gemm_xdl example_grouped_gemm_xdl_bf16)
+        add_example_executable(example_grouped_gemm_xdl_fixed_nk_bias_fp16 grouped_gemm_xdl_fixed_nk_bias_fp16.cpp)
+        add_example_dependencies(example_grouped_gemm example_grouped_gemm_xdl_fixed_nk_bias_fp16)
 
-add_example_executable(example_grouped_gemm_xdl_int8 grouped_gemm_xdl_int8.cpp)
-add_example_dependencies(example_grouped_gemm_xdl example_grouped_gemm_xdl_int8)
+        add_example_executable(example_grouped_gemm_xdl_bf16 grouped_gemm_xdl_bf16.cpp)
+        add_example_dependencies(example_grouped_gemm example_grouped_gemm_xdl_bf16)
 
-add_example_executable(example_grouped_gemm_xdl_fixed_nk_fp8 grouped_gemm_xdl_fixed_nk_fp8.cpp)
-add_dependencies(example_grouped_gemm_xdl example_grouped_gemm_xdl_fixed_nk_fp8)
+        add_example_executable(example_grouped_gemm_xdl_int8 grouped_gemm_xdl_int8.cpp)
+        add_example_dependencies(example_grouped_gemm example_grouped_gemm_xdl_int8)
 
-add_example_executable(example_grouped_gemm_multiple_d_splitk_xdl_fp16 grouped_gemm_multiple_d_splitk_xdl_fp16.cpp)
-add_dependencies(example_grouped_gemm_xdl example_grouped_gemm_multiple_d_splitk_xdl_fp16)
+        add_example_executable(example_grouped_gemm_xdl_fixed_nk_fp8 grouped_gemm_xdl_fixed_nk_fp8.cpp)
+        add_dependencies(example_grouped_gemm example_grouped_gemm_xdl_fixed_nk_fp8)
 
+        add_example_executable(example_grouped_gemm_multiple_d_splitk_xdl_fp16 grouped_gemm_multiple_d_splitk_xdl_fp16.cpp)
+        add_dependencies(example_grouped_gemm example_grouped_gemm_multiple_d_splitk_xdl_fp16)
         
-if(USE_BITINT_EXTENSION_INT4)
+        if(USE_BITINT_EXTENSION_INT4)
             add_example_executable(example_grouped_gemm_xdl_int4 grouped_gemm_xdl_int4.cpp)
-    add_example_dependencies(example_grouped_gemm_xdl example_grouped_gemm_xdl_int4)
-endif()
+            add_example_dependencies(example_grouped_gemm example_grouped_gemm_xdl_int4)
+        endif()
+
+        set(target 1)
+    endif()
+endforeach()
+

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

@@ -72,7 +72,7 @@ __global__ void
                                           const CDEElementwiseOperation cde_element_op)
 {
 #if(!defined(__HIP_DEVICE_COMPILE__) || defined(__gfx908__) || defined(__gfx90a__) || \
-    defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__))
+    defined(__gfx94__))
 
     constexpr index_t shared_size = GridwiseGemm::GetSharedMemoryNumberOfByte();
     __shared__ uint8_t p_shared[shared_size];