diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh
index a328a2a82b..dc472cbecd 100644
--- a/ggml/src/ggml-cuda/common.cuh
+++ b/ggml/src/ggml-cuda/common.cuh
@@ -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;
     }
diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
index e25be3592f..ed2036b9d3 100644
--- a/ggml/src/ggml-cuda/ggml-cuda.cu
+++ b/ggml/src/ggml-cuda/ggml-cuda.cu
@@ -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)