Public/llama.cpp
1
0
Fork 0
mirror of https://github.com/ggml-org/llama.cpp.git synced 2026-09-19 01:04:55 +02:00
Code Issues Packages Projects Releases Wiki Activity

Fix: Perf regression on ada; excludes ada and below from PDL launches

Browse Source
This commit is contained in:
aendk
2026-05-15 16:31:07 +02:00
parent 4346c541a8
commit 378e8e7aae
2 changed files with 18 additions and 7 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-cuda/common.cuh
+4 -7
View File
@@ -224,6 +224,8 @@ struct ggml_cuda_pdl_config {
ggml_cuda_pdl_config& operator=(ggml_cuda_pdl_config&&) = delete;
};
bool ggml_cuda_pdl_disabled();
#endif //defined(GGML_CUDA_USE_PDL)
@@ -231,14 +233,9 @@ template<typename Kernel, typename... Args>
static __inline__ void ggml_cuda_kernel_launch(Kernel kernel, const ggml_cuda_kernel_launch_params & launch_params, Args&&... args) {
#if defined(GGML_CUDA_USE_PDL)
// PDL is on by default. To disable, set GGML_CUDA_DISABLE_PDL=1
static const bool disable_pdl = []() {
const char * env = getenv("GGML_CUDA_DISABLE_PDL");
return env != nullptr && std::atoi(env) != 0;
}();
if (!disable_pdl) {
if (!ggml_cuda_pdl_disabled()) {
auto pdl_cfg = ggml_cuda_pdl_config(launch_params);
CUDA_CHECK(cudaLaunchKernelEx(&pdl_cfg.cfg, kernel, std::forward<Args>(args)... ));
return;
}
ggml/src/ggml-cuda/ggml-cuda.cu
+14
View File
@@ -353,6 +353,20 @@ const ggml_cuda_device_info & ggml_cuda_info() {
return info;
}
#if defined(GGML_CUDA_USE_PDL)
bool ggml_cuda_pdl_disabled() {
static const bool env_disable_pdl = []() {
const char * env = getenv("GGML_CUDA_DISABLE_PDL");
return env != nullptr && std::atoi(env) != 0;
}();
if (env_disable_pdl) {
return true;
}
// Runtime GPU check due to perf regressions on Ada with PDL.
return ggml_cuda_info().devices[ggml_cuda_get_device()].cc < GGML_CUDA_CC_HOPPER;
}
#endif // defined(GGML_CUDA_USE_PDL)
// #define DEBUG_CUDA_MALLOC
// buffer pool for cuda (legacy)