diff --git a/ggml/src/ggml-cuda/fattn-tile.cuh b/ggml/src/ggml-cuda/fattn-tile.cuh index be4f3281a5b8..0f7efdef01db 100644 --- a/ggml/src/ggml-cuda/fattn-tile.cuh +++ b/ggml/src/ggml-cuda/fattn-tile.cuh @@ -68,6 +68,7 @@ static constexpr __host__ __device__ uint32_t ggml_cuda_fattn_tile_get_config_nv GGML_CUDA_FATTN_TILE_CONFIG_CASE(256, 256, 16, 256, 2, 64, 64) GGML_CUDA_FATTN_TILE_CONFIG_CASE(256, 256, 32, 256, 2, 64, 64) + GGML_CUDA_FATTN_TILE_CONFIG_CASE(512, 512, 2, 64, 2, 32, 64) GGML_CUDA_FATTN_TILE_CONFIG_CASE(512, 512, 4, 128, 2, 64, 64) GGML_CUDA_FATTN_TILE_CONFIG_CASE(512, 512, 8, 256, 2, 64, 64) GGML_CUDA_FATTN_TILE_CONFIG_CASE(512, 512, 16, 256, 2, 64, 64) @@ -132,6 +133,7 @@ static constexpr __host__ __device__ uint32_t ggml_cuda_fattn_tile_get_config_nv GGML_CUDA_FATTN_TILE_CONFIG_CASE(256, 256, 16, 256, 2, 32, 128) GGML_CUDA_FATTN_TILE_CONFIG_CASE(256, 256, 32, 256, 2, 32, 64) + GGML_CUDA_FATTN_TILE_CONFIG_CASE(512, 512, 2, 64, 2, 32, 64) GGML_CUDA_FATTN_TILE_CONFIG_CASE(512, 512, 4, 128, 2, 32, 64) GGML_CUDA_FATTN_TILE_CONFIG_CASE(512, 512, 8, 256, 2, 32, 64) GGML_CUDA_FATTN_TILE_CONFIG_CASE(512, 512, 16, 256, 2, 32, 64) @@ -203,6 +205,7 @@ static constexpr __host__ __device__ uint32_t ggml_cuda_fattn_tile_get_config_am GGML_CUDA_FATTN_TILE_CONFIG_CASE(256, 256, 16, 256, 2, 32, 128) GGML_CUDA_FATTN_TILE_CONFIG_CASE(256, 256, 32, 256, 2, 32, 128) + GGML_CUDA_FATTN_TILE_CONFIG_CASE(512, 512, 2, 64, 2, 32, 64) GGML_CUDA_FATTN_TILE_CONFIG_CASE(512, 512, 4, 128, 2, 64, 64) GGML_CUDA_FATTN_TILE_CONFIG_CASE(512, 512, 8, 256, 2, 64, 64) GGML_CUDA_FATTN_TILE_CONFIG_CASE(512, 512, 16, 256, 2, 64, 64) @@ -277,6 +280,7 @@ static constexpr __host__ __device__ uint32_t ggml_cuda_fattn_tile_get_config_am GGML_CUDA_FATTN_TILE_CONFIG_CASE(256, 256, 16, 256, 5, 32, 256) GGML_CUDA_FATTN_TILE_CONFIG_CASE(256, 256, 32, 256, 3, 64, 128) + GGML_CUDA_FATTN_TILE_CONFIG_CASE(512, 512, 2, 64, 2, 32, 64) GGML_CUDA_FATTN_TILE_CONFIG_CASE(512, 512, 4, 128, 2, 64, 64) GGML_CUDA_FATTN_TILE_CONFIG_CASE(512, 512, 8, 256, 2, 64, 64) GGML_CUDA_FATTN_TILE_CONFIG_CASE(512, 512, 16, 256, 4, 64, 64) @@ -1259,6 +1263,16 @@ static void launch_fattn_tile_switch_ncols2(ggml_backend_cuda_context & ctx, ggm launch_fattn_tile_switch_ncols1(ctx, dst); return; } + + // DV > 256 (e.g. DKQ=DV=512, head_dim=512 models): extend GQA fallback to ncols2=2/1. + // Without this, gqa_ratio not divisible by 4 (e.g. ratio=2) reaches GGML_ABORT. + if (use_gqa_opt && gqa_ratio % 2 == 0) { + launch_fattn_tile_switch_ncols1(ctx, dst); + return; + } + + launch_fattn_tile_switch_ncols1(ctx, dst); + return; } GGML_ABORT("fatal error"); } diff --git a/ggml/src/ggml-hip/CMakeLists.txt b/ggml/src/ggml-hip/CMakeLists.txt index f40b1d754f5a..7c70f222e4c6 100644 --- a/ggml/src/ggml-hip/CMakeLists.txt +++ b/ggml/src/ggml-hip/CMakeLists.txt @@ -85,14 +85,17 @@ else() ../ggml-cuda/template-instances/fattn-vec-instance-turbo3_0-turbo3_0.cu ../ggml-cuda/template-instances/fattn-vec-instance-turbo3_0-q8_0.cu ../ggml-cuda/template-instances/fattn-vec-instance-q8_0-turbo3_0.cu + ../ggml-cuda/template-instances/fattn-vec-instance-f16-turbo3_0.cu ../ggml-cuda/template-instances/fattn-vec-instance-turbo2_0-turbo2_0.cu ../ggml-cuda/template-instances/fattn-vec-instance-turbo2_0-q8_0.cu ../ggml-cuda/template-instances/fattn-vec-instance-q8_0-turbo2_0.cu + ../ggml-cuda/template-instances/fattn-vec-instance-f16-turbo2_0.cu ../ggml-cuda/template-instances/fattn-vec-instance-turbo3_0-turbo2_0.cu ../ggml-cuda/template-instances/fattn-vec-instance-turbo2_0-turbo3_0.cu ../ggml-cuda/template-instances/fattn-vec-instance-turbo4_0-turbo4_0.cu ../ggml-cuda/template-instances/fattn-vec-instance-turbo4_0-q8_0.cu ../ggml-cuda/template-instances/fattn-vec-instance-q8_0-turbo4_0.cu + ../ggml-cuda/template-instances/fattn-vec-instance-f16-turbo4_0.cu ../ggml-cuda/template-instances/fattn-vec-instance-turbo4_0-turbo3_0.cu ../ggml-cuda/template-instances/fattn-vec-instance-turbo3_0-turbo4_0.cu ../ggml-cuda/template-instances/fattn-vec-instance-turbo4_0-turbo2_0.cu