diff --git a/3rdparty/QoLA b/3rdparty/QoLA
index 47e0f5501a..0d6665b6d8 160000
--- a/3rdparty/QoLA
+++ b/3rdparty/QoLA
@@ -1 +1 @@
-Subproject commit 47e0f5501af51521c8a911fbbefa4e24b92dd7b6
+Subproject commit 0d6665b6d87feeafe4b782dcc736c5e755dc5b34
diff --git a/transformer_engine/common/ck_fused_attn/include/ck_fused_attn/ck_fused_attn.hpp b/transformer_engine/common/ck_fused_attn/include/ck_fused_attn/ck_fused_attn.hpp
index 127d75b4ca..150c8bfb32 100644
--- a/transformer_engine/common/ck_fused_attn/include/ck_fused_attn/ck_fused_attn.hpp
+++ b/transformer_engine/common/ck_fused_attn/include/ck_fused_attn/ck_fused_attn.hpp
@@ -143,6 +143,12 @@ struct CkAttnBwdArgs : CKAttnCommonArgs {
 hipError_t ck_attn_fwd(const CKAttnFwdArgs& args, hipStream_t stream);
 hipError_t ck_attn_bwd(const CkAttnBwdArgs& args, hipStream_t stream);
 
+// Probe whether AITER's v3 (asm) path will run for the given config, without
+// launching a kernel (backed by AITER's v3_api_check dry-run). Returns true iff
+// the v3 path is selected; false means the CK v2 path (or no support) would run.
+bool ck_attn_fwd_uses_v3(const CKAttnFwdArgs& args);
+bool ck_attn_bwd_uses_v3(const CkAttnBwdArgs& args);
+
 }//namespace ck_fused_attn
 #endif // CK_FUSED_ATTN_H
 
diff --git a/transformer_engine/common/ck_fused_attn/src/ck_fused_attn_bwd.cpp b/transformer_engine/common/ck_fused_attn/src/ck_fused_attn_bwd.cpp
index b0ed9ae6e5..a7425be988 100644
--- a/transformer_engine/common/ck_fused_attn/src/ck_fused_attn_bwd.cpp
+++ b/transformer_engine/common/ck_fused_attn/src/ck_fused_attn_bwd.cpp
@@ -441,21 +441,24 @@ void dump_bwd_timings(const char* dump_path, float average_runtime){
   file << average_runtime << "\n";
 }
 
-hipError_t ck_attn_bwd(const CkAttnBwdArgs& args, hipStream_t stream){
+struct BwdFmhaArgs {
+  aiter::mha_bwd_args fmha_args;
+  bool is_mqa_gqa;
+  bool has_dbias;
+  BiasShape bias_shape;
+};
+
+// Populate the AITER mha_bwd_args from TE's CkAttnBwdArgs. Shared by ck_attn_bwd
+// (real launch) and ck_attn_bwd_uses_v3 (v3 availability probe) so the probe can
+// never disagree with the launch. v3_api_check is left false here; callers flip it.
+// The stream-dependent max_seqlen override (NVTE_CK_RUNTIME_MAX_SEQLEN) is applied
+// by ck_attn_bwd after this returns; it does not affect v3 kernel selection.
+BwdFmhaArgs build_bwd_fmha_args(const CkAttnBwdArgs& args){
 
   bool has_dropout = (args.dropout_probability > 0.f);
   bool has_dbias = args.dbias_ptr != nullptr;
   bool is_mqa_gqa = (args.h > args.hg);
 
-  bool ck_log_config = false;
-  if (const char* env_p = std::getenv("CK_FUSED_ATTN_LOG_CONFIG") ) {
-    if (env_p != nullptr && std::string(env_p) == "1")
-      ck_log_config = true;
-  }
-  const char* dump_path = std::getenv("NVTE_DUMP_AITER_RT");
-  // print kernel name on verbose mode
-  ck_tile::stream_config stream_config{stream, dump_path!=nullptr, get_ck_log_stream() != nullptr};
-
   bias_enum bias_type = bias_enum::no_bias;
   BiasShape bias_shape = BiasShape::k11SS;
   if (!args.is_group_mode()) {
@@ -578,6 +581,39 @@ hipError_t ck_attn_bwd(const CkAttnBwdArgs& args, hipStream_t stream){
   fmha_args.p_undrop = 1.0 - args.dropout_probability;
   fmha_args.drop_seed_offset = std::pair<const void*, const void*>{args.philox_seed_ptr, args.philox_offset_ptr};
 
+  return {fmha_args, is_mqa_gqa, has_dbias, bias_shape};
+}
+
+// Probe whether AITER's v3 (asm) backward path will run for this config, without
+// launching any kernel. Builds the same args as ck_attn_bwd and relies on AITER's
+// v3_api_check dry-run (returns 1 when v3 is available, -1 otherwise).
+bool ck_attn_bwd_uses_v3(const CkAttnBwdArgs& args){
+  aiter::mha_bwd_args fmha_args = build_bwd_fmha_args(args).fmha_args;
+  fmha_args.v3_api_check = true;
+  // No kernel is launched in check mode, so the stream/log flags are irrelevant.
+  ck_tile::stream_config stream_config{nullptr, false, false};
+  return QOLA_NS(mha_bwd)(fmha_args, stream_config) == 1;
+}
+
+hipError_t ck_attn_bwd(const CkAttnBwdArgs& args, hipStream_t stream){
+
+  // build_bwd_fmha_args also derives the flags the post-dispatch MQA/GQA and dbias
+  // reductions below need, so we unpack them here rather than recomputing.
+  BwdFmhaArgs built = build_bwd_fmha_args(args);
+  aiter::mha_bwd_args& fmha_args = built.fmha_args;
+  const bool is_mqa_gqa = built.is_mqa_gqa;
+  const bool has_dbias = built.has_dbias;
+  const BiasShape bias_shape = built.bias_shape;
+
+  bool ck_log_config = false;
+  if (const char* env_p = std::getenv("CK_FUSED_ATTN_LOG_CONFIG") ) {
+    if (env_p != nullptr && std::string(env_p) == "1")
+      ck_log_config = true;
+  }
+  const char* dump_path = std::getenv("NVTE_DUMP_AITER_RT");
+  // print kernel name on verbose mode
+  ck_tile::stream_config stream_config{stream, dump_path!=nullptr, get_ck_log_stream() != nullptr};
+
   // modify the max_seqlen_q for better performance in 0-length cases
   // lse_workspace_ptr used as buffer
   if(const char* env_p = std::getenv("NVTE_CK_RUNTIME_MAX_SEQLEN")) {
diff --git a/transformer_engine/common/ck_fused_attn/src/ck_fused_attn_fwd.cpp b/transformer_engine/common/ck_fused_attn/src/ck_fused_attn_fwd.cpp
index 0f407230c0..88e3bac91a 100644
--- a/transformer_engine/common/ck_fused_attn/src/ck_fused_attn_fwd.cpp
+++ b/transformer_engine/common/ck_fused_attn/src/ck_fused_attn_fwd.cpp
@@ -99,18 +99,12 @@ void dump_fwd_timings(const char* dump_path, float average_runtime){
   file << average_runtime << "\n";
 }
 
-hipError_t ck_attn_fwd(const CKAttnFwdArgs& args, hipStream_t stream){
-
-  bool has_dropout = (args.is_training && args.dropout_probability > 0.f);
-
-  bool ck_log_config = false;
-  if (const char* env_p = std::getenv("CK_FUSED_ATTN_LOG_CONFIG") ) {
-    if (env_p != nullptr && std::string(env_p) == "1")
-      ck_log_config = true;
-  }
-  const char* dump_path = std::getenv("NVTE_DUMP_AITER_RT");
-  // print kernel name on verbose mode
-  ck_tile::stream_config stream_config{stream, dump_path!=nullptr, get_ck_log_stream() != nullptr};
+// Populate the AITER mha_fwd_args from TE's CKAttnFwdArgs. Shared by ck_attn_fwd
+// (real launch) and ck_attn_fwd_uses_v3 (v3 availability probe) so the probe can
+// never disagree with the launch. v3_api_check is left false here; callers flip it.
+// The stream-dependent max_seqlen override (NVTE_CK_RUNTIME_MAX_SEQLEN) is applied
+// by ck_attn_fwd after this returns; it does not affect v3 kernel selection.
+aiter::mha_fwd_args build_fwd_fmha_args(const CKAttnFwdArgs& args){
 
   bias_enum bias_type = bias_enum::no_bias;
   BiasShape bias_shape = BiasShape::k11SS;
@@ -216,6 +210,35 @@ hipError_t ck_attn_fwd(const CKAttnFwdArgs& args, hipStream_t stream){
   fmha_args.block_scale_size_q  = 0;
   fmha_args.block_scale_size_kv = 0;
 
+  return fmha_args;
+}
+
+// Probe whether AITER's v3 (asm) forward path will run for this config, without
+// launching any kernel. Builds the same args as ck_attn_fwd and relies on AITER's
+// v3_api_check dry-run (returns 1 when v3 is available, -1 otherwise).
+bool ck_attn_fwd_uses_v3(const CKAttnFwdArgs& args){
+  aiter::mha_fwd_args fmha_args = build_fwd_fmha_args(args);
+  fmha_args.v3_api_check = true;
+  // No kernel is launched in check mode, so the stream/log flags are irrelevant.
+  ck_tile::stream_config stream_config{nullptr, false, false};
+  return QOLA_NS(mha_fwd)(fmha_args, stream_config) == 1;
+}
+
+hipError_t ck_attn_fwd(const CKAttnFwdArgs& args, hipStream_t stream){
+
+  bool has_dropout = (args.is_training && args.dropout_probability > 0.f);
+
+  bool ck_log_config = false;
+  if (const char* env_p = std::getenv("CK_FUSED_ATTN_LOG_CONFIG") ) {
+    if (env_p != nullptr && std::string(env_p) == "1")
+      ck_log_config = true;
+  }
+  const char* dump_path = std::getenv("NVTE_DUMP_AITER_RT");
+  // print kernel name on verbose mode
+  ck_tile::stream_config stream_config{stream, dump_path!=nullptr, get_ck_log_stream() != nullptr};
+
+  aiter::mha_fwd_args fmha_args = build_fwd_fmha_args(args);
+
   if(const char* env_p = std::getenv("NVTE_CK_RUNTIME_MAX_SEQLEN")){
     if(args.is_group_mode() && std::string(env_p) == "1"){
       if(ck_log_config){