nicolas
/
ik_llama_opt
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Adjust CUDA FA kernel parameters for head size 512 on Turing (#1942)

This commit is contained in:
Kawrakow 2026-06-10 07:49:21 +02:00 committed by GitHub
parent 2a1148384c
commit e6f8112f3b
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
2 changed files with 60 additions and 14 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
examples/server/server.cpp
View File

@ -1141,8 +1141,8 @@ int main(int argc, char ** argv) {
arr.push_back(res->to_json());
}
// if single request, return single object instead of array
res_ok(res, arr.size() == 1 ? arr[0] : arr);
}
res_ok(res, arr.size() == 1 ? arr[0] : arr);
}
}
else {
// in streaming mode, the first error must be treated as non-stream response

70
ggml/src/ggml-cuda/fattn-new-mma.cu
View File

@ -332,35 +332,81 @@ struct fattn_mma_f16_config<192, 192> {
}
};
//template <>
//struct fattn_mma_f16_config<512, 512> {
// static constexpr int nbatch_fa = 64;
// static constexpr int nwarps_max = 4;
// static constexpr bool Q_in_reg = false;
// static constexpr int nstages_target = 1;
//
// static int get_nbatch_K2_host([[maybe_unused]] const int cc, [[maybe_unused]] const int ncols) {
// return 256;
// }
//
// static constexpr __device__ int get_nbatch_K2_device([[maybe_unused]] int ncols) {
// return 256;
// }
//
// static int get_nbatch_V2_host([[maybe_unused]] const int cc, [[maybe_unused]] const int ncols) {
// return 256;
// }
//
// static constexpr __device__ int get_nbatch_V2_device([[maybe_unused]] int ncols) {
// return 256;
// }
//
// static int get_nbatch_combine_host(const int /*cc*/, const int /*ncols*/) {
// return 256;
// }
//
// static constexpr __device__ int get_nbatch_combine_device(int /*ncols*/) {
// return 256;
// }
//};
template <>
struct fattn_mma_f16_config<512, 512> {
static constexpr int nbatch_fa = 64;
static constexpr int nwarps_max = 4;
static constexpr int nbatch_fa = 32;
static constexpr int nwarps_max = 8;
static constexpr bool Q_in_reg = false;
static constexpr int nstages_target = 1;
static int get_nbatch_K2_host([[maybe_unused]] const int cc, [[maybe_unused]] const int ncols) {
return 256;
static int get_nbatch_K2_host(const int cc, const int ncols) {
if (ggml_cuda_highest_compiled_arch(cc) == CC_TURING) {
return ncols <= 16 ? 96 : 128;
}
return ncols <= 16 ? 256 : 128;
}
static constexpr __device__ int get_nbatch_K2_device([[maybe_unused]] int ncols) {
return 256;
static constexpr __device__ int get_nbatch_K2_device(int ncols) {
#if __CUDA_ARCH__ == CC_TURING
return ncols <= 16 ? 96 : 128;
#else
return ncols <= 16 ? 256 : 128;
#endif // __CUDA_ARCH__ == CC_TURING
}
static int get_nbatch_V2_host([[maybe_unused]] const int cc, [[maybe_unused]] const int ncols) {
return 256;
static int get_nbatch_V2_host(const int cc, const int ncols) {
if (ggml_cuda_highest_compiled_arch(cc) == CC_TURING) {
return ncols <= 16 ? 64 : 128;
}
return ncols <= 16 ? 256 : 128;
}
static constexpr __device__ int get_nbatch_V2_device([[maybe_unused]] int ncols) {
return 256;
static constexpr __device__ int get_nbatch_V2_device(int ncols) {
#if __CUDA_ARCH__ == CC_TURING
return ncols <= 16 ? 64 : 128;
#else
return ncols <= 16 ? 256 : 128;
#endif // __CUDA_ARCH__ == CC_TURING
}
static int get_nbatch_combine_host(const int /*cc*/, const int /*ncols*/) {
return 256;
return 128;
}
static constexpr __device__ int get_nbatch_combine_device(int /*ncols*/) {
return 256;
return 128;
}
};