@@ -89,10 +89,10 @@ static constexpr __host__ __device__ fattn_mma_config ggml_cuda_fattn_mma_get_co
8989}
9090
9191static constexpr __host__ __device__ fattn_mma_config ggml_cuda_fattn_mma_get_config_volta (const int DKQ, const int DV, const int ncols) {
92- // GGML_CUDA_FATTN_MMA_CONFIG_CASE(576, 512, 8, 64, 4, 32, 288, 256, 64, 1, false);
93- // GGML_CUDA_FATTN_MMA_CONFIG_CASE(576, 512, 16, 64, 4, 32, 288, 256, 64, 1, false);
94- // GGML_CUDA_FATTN_MMA_CONFIG_CASE(576, 512, 32, 128, 2, 32, 160, 128, 64, 1, false);
95- // GGML_CUDA_FATTN_MMA_CONFIG_CASE(576, 512, 64, 256, 1, 32, 160, 128, 64, 1, false);
92+ GGML_CUDA_FATTN_MMA_CONFIG_CASE (576 , 512 , 8 , 64 , 4 , 32 , 288 , 256 , 64 , 1 , false );
93+ GGML_CUDA_FATTN_MMA_CONFIG_CASE (576 , 512 , 16 , 64 , 4 , 32 , 288 , 256 , 64 , 1 , false );
94+ GGML_CUDA_FATTN_MMA_CONFIG_CASE (576 , 512 , 32 , 128 , 2 , 32 , 160 , 128 , 64 , 1 , false );
95+ GGML_CUDA_FATTN_MMA_CONFIG_CASE (576 , 512 , 64 , 256 , 1 , 32 , 160 , 128 , 64 , 1 , false );
9696
9797 // TODO tune specifically for Volta
9898 return ggml_cuda_fattn_mma_get_config_ampere (DKQ, DV, ncols);
0 commit comments