support group_size=24 (#7636)

Author: chang-wenbin

custom_ops/gpu_ops/append_attn/template_config.json

@@ -17,7 +17,7 @@
       "IsDynamicC8"
     ],
     "dispatch_params": {
-      "GROUP_SIZE": [1, 2, 4, 5, 6, 7, 8, 12, 14, 16],
+      "GROUP_SIZE": [1, 2, 4, 5, 6, 7, 8, 12, 14, 16, 24],
       "HEAD_DIM": [128],
       "BLOCK_SIZE": [64],
       "CAUSAL": [0, 1],
@@ -54,7 +54,7 @@
       "ENABLE_PREFILL"
     ],
     "dispatch_params": {
-      "GROUP_SIZE": [1, 2, 4, 5, 6, 7, 8, 12, 14, 16],
+      "GROUP_SIZE": [1, 2, 4, 5, 6, 7, 8, 12, 14, 16, 24],
       "HEAD_DIM": [128],
       "BLOCK_SIZE": [64],
       "CAUSAL": [0, 1],
@@ -89,7 +89,7 @@
       "ENABLE_PREFILL"
     ],
     "dispatch_params": {
-      "GROUP_SIZE": [1, 2, 4, 5, 6, 7, 8, 12, 14, 16],
+      "GROUP_SIZE": [1, 2, 4, 5, 6, 7, 8, 12, 14, 16, 24],
       "HEAD_DIM": [64,128],
       "BLOCK_SIZE": [64],
       "CAUSAL": [0, 1],

custom_ops/gpu_ops/append_attn/utils.cuh

@@ -445,6 +445,9 @@ __forceinline__ __host__ __device__ void vec_cast<nv_bfloat16, float>(
   } else if (group_size == 16) {                             \
     constexpr size_t GROUP_SIZE = 16;                        \
     __VA_ARGS__                                              \
+  } else if (group_size == 24) {                             \
+    constexpr size_t GROUP_SIZE = 24;                        \
+    __VA_ARGS__                                              \
   } else {                                                   \
     PD_THROW("not support the group_size", group_size);      \
   }