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 3f51b96b6..0bd062f6b 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
@@ -1,5 +1,5 @@
 /*************************************************************************
- * Copyright (c) 2024-2025, Advanced Micro Devices, Inc. All rights reserved.
+ * Copyright (c) 2024-2026, Advanced Micro Devices, Inc. All rights reserved.
  *
  * License for AMD contributions = MIT. See LICENSE for more information
  ************************************************************************/
@@ -16,6 +16,34 @@
 
 namespace ck_fused_attn{
 
+// We want to cache and reuse the log stream so we use thread_local here.
+namespace {
+std::ostream* get_bwd_log_stream() {
+  thread_local std::ofstream log_file;
+  thread_local bool attempted = false;
+  thread_local bool opened = false;
+  thread_local bool requested = false;
+  thread_local std::string log_dir_str;
+  if (!attempted) {
+    attempted = true;
+    if (const char* env_p = std::getenv("CK_FUSED_ATTN_LOG_CONFIG")) {
+      log_dir_str = std::string(env_p);
+      requested = !log_dir_str.empty() && log_dir_str != "0";
+    }
+    if (requested) {
+      opened = open_ck_fused_attn_log_file(log_file, "ck_fused_attn_bwd", log_dir_str);
+    }
+  }
+  if (!requested) {
+    return nullptr;
+  }
+  if (!opened) {
+    return &std::cout;
+  }
+  return &log_file;
+}
+}  // namespace
+
 // TODO: unify with binary search in TE/common/fused_attn(rocm)/util
 // no device std::upper_bound
 // in an increasing array with given size len, search for the index that:
@@ -346,110 +374,104 @@ void log_bwd_config(const char* func_name,
                     const bool is_v3_atomic_fp32,
                     const int how_v3_bf16_cvt,
                     const fmha_bwd_args& fmha_args){
-
-  bool ck_fused_attn_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_fused_attn_log_config = true;
-  }
-  if (ck_fused_attn_log_config) {
-    std::cout<<std::endl<<func_name<<std::endl;
+  if (auto* log_file = get_bwd_log_stream()) {
+    *log_file << "\n" << func_name << "\n";
 
     // fmha_traits debug
-    std::cout<<std::endl<<"fmha_traits: "<<std::endl;
-    std::cout<<"hdim_q: "<<fmha_args.hdim_q<<std::endl;
-    std::cout<<"hdim_v: "<<fmha_args.hdim_v<<std::endl;
-    std::cout<<"data_type: "<<data_type_str<<std::endl;
-    std::cout<<"is_group_mode: "<<is_group_mode<<std::endl;
-    std::cout<<"mask_type: "<<static_cast<std::underlying_type<mask_enum>::type>(mask_type)<<std::endl;
-    std::cout<<"bias_type: "<<static_cast<std::underlying_type<bias_enum>::type>(bias_type)<<std::endl;
-    std::cout<<"has_dbias: "<<has_dbias<<std::endl;
-    std::cout<<"has_dropout: "<<has_dropout<<std::endl;
-    std::cout<<"is_store_randval: "<<is_store_randval<<std::endl;
-    std::cout<<"is_deterministic: "<<is_deterministic<<std::endl;
-    std::cout<<"uses_bwd_v3: "<<uses_bwd_v3<<std::endl;
-    std::cout<<"is_v3_atomic_fp32: "<<is_v3_atomic_fp32<<std::endl;
-    std::cout<<"how_v3_bf16_cvt: "<<how_v3_bf16_cvt<<std::endl;
+    *log_file << "\n" << "fmha_traits: " << "\n";
+    *log_file << "hdim_q: " << fmha_args.hdim_q << "\n";
+    *log_file << "hdim_v: " << fmha_args.hdim_v << "\n";
+    *log_file << "data_type: " << data_type_str << "\n";
+    *log_file << "is_group_mode: " << is_group_mode << "\n";
+    *log_file << "mask_type: " << static_cast<std::underlying_type<mask_enum>::type>(mask_type) << "\n";
+    *log_file << "bias_type: " << static_cast<std::underlying_type<bias_enum>::type>(bias_type) << "\n";
+    *log_file << "has_dbias: " << has_dbias << "\n";
+    *log_file << "has_dropout: " << has_dropout << "\n";
+    *log_file << "is_store_randval: " << is_store_randval << "\n";
+    *log_file << "is_deterministic: " << is_deterministic << "\n";
+    *log_file << "uses_bwd_v3: " << uses_bwd_v3 << "\n";
+    *log_file << "is_v3_atomic_fp32: " << is_v3_atomic_fp32 << "\n";
+    *log_file << "how_v3_bf16_cvt: " << how_v3_bf16_cvt << "\n";
 
     // fmha_args debug
-    std::cout<<std::endl<<"fmha_args: "<<std::endl;
-    std::cout<<"q_ptr: "<<fmha_args.q_ptr<<std::endl;
-    std::cout<<"k_ptr: "<<fmha_args.k_ptr<<std::endl;
-    std::cout<<"v_ptr: "<<fmha_args.v_ptr<<std::endl;
-    std::cout<<"bias_ptr: "<<fmha_args.bias_ptr<<std::endl;
-    std::cout<<"o_ptr: "<<fmha_args.o_ptr<<std::endl;
-    std::cout<<"lse_ptr: "<<fmha_args.lse_ptr<<std::endl;
-    std::cout<<"do_ptr: "<<fmha_args.do_ptr<<std::endl;
-    std::cout<<"d_ptr: "<<fmha_args.d_ptr<<std::endl;
-    std::cout<<"rand_val_ptr: "<<fmha_args.rand_val_ptr<<std::endl;
-    std::cout<<"dq_ptr: "<<fmha_args.dq_ptr<<std::endl;
-    std::cout<<"dk_ptr: "<<fmha_args.dk_ptr<<std::endl;
-    std::cout<<"dv_ptr: "<<fmha_args.dv_ptr<<std::endl;
-    std::cout<<"dbias_ptr: "<<fmha_args.dbias_ptr<<std::endl;
-    std::cout<<"dq_acc_ptr: "<<fmha_args.dq_acc_ptr<<std::endl;
+    *log_file << "\n" << "fmha_args: " << "\n";
+    *log_file << "q_ptr: " << fmha_args.q_ptr << "\n";
+    *log_file << "k_ptr: " << fmha_args.k_ptr << "\n";
+    *log_file << "v_ptr: " << fmha_args.v_ptr << "\n";
+    *log_file << "bias_ptr: " << fmha_args.bias_ptr << "\n";
+    *log_file << "o_ptr: " << fmha_args.o_ptr << "\n";
+    *log_file << "lse_ptr: " << fmha_args.lse_ptr << "\n";
+    *log_file << "do_ptr: " << fmha_args.do_ptr << "\n";
+    *log_file << "d_ptr: " << fmha_args.d_ptr << "\n";
+    *log_file << "rand_val_ptr: " << fmha_args.rand_val_ptr << "\n";
+    *log_file << "dq_ptr: " << fmha_args.dq_ptr << "\n";
+    *log_file << "dk_ptr: " << fmha_args.dk_ptr << "\n";
+    *log_file << "dv_ptr: " << fmha_args.dv_ptr << "\n";
+    *log_file << "dbias_ptr: " << fmha_args.dbias_ptr << "\n";
+    *log_file << "dq_acc_ptr: " << fmha_args.dq_acc_ptr << "\n";
 
-    std::cout<<"seqstart_q_ptr: "<<fmha_args.seqstart_q_ptr<<std::endl;
-    std::cout<<"seqstart_k_ptr: "<<fmha_args.seqstart_k_ptr<<std::endl;
-    std::cout<<"seqlen_q_ptr: "<<fmha_args.seqlen_q_ptr<<std::endl;
-    std::cout<<"seqlen_k_ptr: "<<fmha_args.seqlen_k_ptr<<std::endl;
-    std::cout<<"cu_seqlen_q_ptr: "<<fmha_args.cu_seqlen_q_ptr<<std::endl;
-    std::cout<<"cu_seqlen_k_ptr: "<<fmha_args.cu_seqlen_k_ptr<<std::endl;
+    *log_file << "seqstart_q_ptr: " << fmha_args.seqstart_q_ptr << "\n";
+    *log_file << "seqstart_k_ptr: " << fmha_args.seqstart_k_ptr << "\n";
+    *log_file << "seqlen_q_ptr: " << fmha_args.seqlen_q_ptr << "\n";
+    *log_file << "seqlen_k_ptr: " << fmha_args.seqlen_k_ptr << "\n";
+    *log_file << "cu_seqlen_q_ptr: " << fmha_args.cu_seqlen_q_ptr << "\n";
+    *log_file << "cu_seqlen_k_ptr: " << fmha_args.cu_seqlen_k_ptr << "\n";
 
-    std::cout<<"seqlen_q: "<<fmha_args.seqlen_q<<std::endl;
-    std::cout<<"seqlen_k: "<<fmha_args.seqlen_k<<std::endl;
-    std::cout<<"batch: "<<fmha_args.batch<<std::endl;
-    std::cout<<"max_seqlen_q: "<<fmha_args.max_seqlen_q<<std::endl;
-    std::cout<<"max_seqlen_k: "<<fmha_args.max_seqlen_k<<std::endl;
-    std::cout<<"hdim_q: "<<fmha_args.hdim_q<<std::endl;
-    std::cout<<"hdim_v: "<<fmha_args.hdim_v<<std::endl;
-    std::cout<<"nhead_q: "<<fmha_args.nhead_q<<std::endl;
-    std::cout<<"nhead_k: "<<fmha_args.nhead_k<<std::endl;
-    std::cout<<"scale: "<<fmha_args.scale<<std::endl;
-    std::cout<<"stride_q: "<<fmha_args.stride_q<<std::endl;
-    std::cout<<"stride_k: "<<fmha_args.stride_k<<std::endl;
-    std::cout<<"stride_v: "<<fmha_args.stride_v<<std::endl;
-    std::cout<<"stride_bias: "<<fmha_args.stride_bias<<std::endl;
-    std::cout<<"stride_o: "<<fmha_args.stride_o<<std::endl;
-    std::cout<<"stride_randval: "<<fmha_args.stride_randval<<std::endl;
-    std::cout<<"stride_do: "<<fmha_args.stride_do<<std::endl;
-    std::cout<<"stride_dq_acc: "<<fmha_args.stride_dq_acc<<std::endl;
-    std::cout<<"stride_dq: "<<fmha_args.stride_dq<<std::endl;
-    std::cout<<"stride_dk: "<<fmha_args.stride_dk<<std::endl;
-    std::cout<<"stride_dv: "<<fmha_args.stride_dv<<std::endl;
-    std::cout<<"stride_dbias: "<<fmha_args.stride_dbias<<std::endl;
-    std::cout<<"nhead_stride_q: "<<fmha_args.nhead_stride_q<<std::endl;
-    std::cout<<"nhead_stride_k: "<<fmha_args.nhead_stride_k<<std::endl;
-    std::cout<<"nhead_stride_v: "<<fmha_args.nhead_stride_v<<std::endl;
-    std::cout<<"nhead_stride_bias: "<<fmha_args.nhead_stride_bias<<std::endl;
-    std::cout<<"nhead_stride_o: "<<fmha_args.nhead_stride_o<<std::endl;
-    std::cout<<"nhead_stride_randval: "<<fmha_args.nhead_stride_randval<<std::endl;
-    std::cout<<"nhead_stride_do: "<<fmha_args.nhead_stride_do<<std::endl;
-    std::cout<<"nhead_stride_lsed: "<<fmha_args.nhead_stride_lsed<<std::endl;
-    std::cout<<"nhead_stride_dq_acc: "<<fmha_args.nhead_stride_dq_acc<<std::endl;
-    std::cout<<"nhead_stride_dq: "<<fmha_args.nhead_stride_dq<<std::endl;
-    std::cout<<"nhead_stride_dk: "<<fmha_args.nhead_stride_dk<<std::endl;
-    std::cout<<"nhead_stride_dv: "<<fmha_args.nhead_stride_dv<<std::endl;
-    std::cout<<"nhead_stride_dbias: "<<fmha_args.nhead_stride_dbias<<std::endl;
-    std::cout<<"batch_stride_q: "<<fmha_args.batch_stride_q<<std::endl;
-    std::cout<<"batch_stride_k: "<<fmha_args.batch_stride_k<<std::endl;
-    std::cout<<"batch_stride_v: "<<fmha_args.batch_stride_v<<std::endl;
-    std::cout<<"batch_stride_bias: "<<fmha_args.batch_stride_bias<<std::endl;
-    std::cout<<"batch_stride_o: "<<fmha_args.batch_stride_o<<std::endl;
-    std::cout<<"batch_stride_randval: "<<fmha_args.batch_stride_randval<<std::endl;
-    std::cout<<"batch_stride_do: "<<fmha_args.batch_stride_do<<std::endl;
-    std::cout<<"batch_stride_lsed: "<<fmha_args.batch_stride_lsed<<std::endl;
-    std::cout<<"batch_stride_dq_acc: "<<fmha_args.batch_stride_dq_acc<<std::endl;
-    std::cout<<"batch_stride_dq: "<<fmha_args.batch_stride_dq<<std::endl;
-    std::cout<<"batch_stride_dk: "<<fmha_args.batch_stride_dk<<std::endl;
-    std::cout<<"batch_stride_dv: "<<fmha_args.batch_stride_dv<<std::endl;
-    std::cout<<"batch_stride_dbias: "<<fmha_args.batch_stride_dbias<<std::endl;
-    std::cout<<"window_size_left: "<<fmha_args.window_size_left<<std::endl;
-    std::cout<<"window_size_right: "<<fmha_args.window_size_right<<std::endl;
-    std::cout<<"mask_type: "<<fmha_args.mask_type<<std::endl;
-    std::cout<<"p_drop: "<<fmha_args.p_drop<<std::endl;
-    std::cout<<"p_undrop: "<<fmha_args.p_undrop<<std::endl;
-    std::cout<<"dropout_seed_ptr: "<<std::get<0>(std::get<std::pair<const void*, const void*>>(fmha_args.drop_seed_offset))<<std::endl;
-    std::cout<<"dropout_offset_ptr: "<<std::get<1>(std::get<std::pair<const void*, const void*>>(fmha_args.drop_seed_offset))<<std::endl;
+    *log_file << "seqlen_q: " << fmha_args.seqlen_q << "\n";
+    *log_file << "seqlen_k: " << fmha_args.seqlen_k << "\n";
+    *log_file << "batch: " << fmha_args.batch << "\n";
+    *log_file << "max_seqlen_q: " << fmha_args.max_seqlen_q << "\n";
+    *log_file << "max_seqlen_k: " << fmha_args.max_seqlen_k << "\n";
+    *log_file << "hdim_q: " << fmha_args.hdim_q << "\n";
+    *log_file << "hdim_v: " << fmha_args.hdim_v << "\n";
+    *log_file << "nhead_q: " << fmha_args.nhead_q << "\n";
+    *log_file << "nhead_k: " << fmha_args.nhead_k << "\n";
+    *log_file << "scale: " << fmha_args.scale << "\n";
+    *log_file << "stride_q: " << fmha_args.stride_q << "\n";
+    *log_file << "stride_k: " << fmha_args.stride_k << "\n";
+    *log_file << "stride_v: " << fmha_args.stride_v << "\n";
+    *log_file << "stride_bias: " << fmha_args.stride_bias << "\n";
+    *log_file << "stride_o: " << fmha_args.stride_o << "\n";
+    *log_file << "stride_randval: " << fmha_args.stride_randval << "\n";
+    *log_file << "stride_do: " << fmha_args.stride_do << "\n";
+    *log_file << "stride_dq_acc: " << fmha_args.stride_dq_acc << "\n";
+    *log_file << "stride_dq: " << fmha_args.stride_dq << "\n";
+    *log_file << "stride_dk: " << fmha_args.stride_dk << "\n";
+    *log_file << "stride_dv: " << fmha_args.stride_dv << "\n";
+    *log_file << "stride_dbias: " << fmha_args.stride_dbias << "\n";
+    *log_file << "nhead_stride_q: " << fmha_args.nhead_stride_q << "\n";
+    *log_file << "nhead_stride_k: " << fmha_args.nhead_stride_k << "\n";
+    *log_file << "nhead_stride_v: " << fmha_args.nhead_stride_v << "\n";
+    *log_file << "nhead_stride_bias: " << fmha_args.nhead_stride_bias << "\n";
+    *log_file << "nhead_stride_o: " << fmha_args.nhead_stride_o << "\n";
+    *log_file << "nhead_stride_randval: " << fmha_args.nhead_stride_randval << "\n";
+    *log_file << "nhead_stride_do: " << fmha_args.nhead_stride_do << "\n";
+    *log_file << "nhead_stride_lsed: " << fmha_args.nhead_stride_lsed << "\n";
+    *log_file << "nhead_stride_dq_acc: " << fmha_args.nhead_stride_dq_acc << "\n";
+    *log_file << "nhead_stride_dq: " << fmha_args.nhead_stride_dq << "\n";
+    *log_file << "nhead_stride_dk: " << fmha_args.nhead_stride_dk << "\n";
+    *log_file << "nhead_stride_dv: " << fmha_args.nhead_stride_dv << "\n";
+    *log_file << "nhead_stride_dbias: " << fmha_args.nhead_stride_dbias << "\n";
+    *log_file << "batch_stride_q: " << fmha_args.batch_stride_q << "\n";
+    *log_file << "batch_stride_k: " << fmha_args.batch_stride_k << "\n";
+    *log_file << "batch_stride_v: " << fmha_args.batch_stride_v << "\n";
+    *log_file << "batch_stride_bias: " << fmha_args.batch_stride_bias << "\n";
+    *log_file << "batch_stride_o: " << fmha_args.batch_stride_o << "\n";
+    *log_file << "batch_stride_randval: " << fmha_args.batch_stride_randval << "\n";
+    *log_file << "batch_stride_do: " << fmha_args.batch_stride_do << "\n";
+    *log_file << "batch_stride_lsed: " << fmha_args.batch_stride_lsed << "\n";
+    *log_file << "batch_stride_dq_acc: " << fmha_args.batch_stride_dq_acc << "\n";
+    *log_file << "batch_stride_dq: " << fmha_args.batch_stride_dq << "\n";
+    *log_file << "batch_stride_dk: " << fmha_args.batch_stride_dk << "\n";
+    *log_file << "batch_stride_dv: " << fmha_args.batch_stride_dv << "\n";
+    *log_file << "batch_stride_dbias: " << fmha_args.batch_stride_dbias << "\n";
+    *log_file << "window_size_left: " << fmha_args.window_size_left << "\n";
+    *log_file << "window_size_right: " << fmha_args.window_size_right << "\n";
+    *log_file << "mask_type: " << fmha_args.mask_type << "\n";
+    *log_file << "p_drop: " << fmha_args.p_drop << "\n";
+    *log_file << "p_undrop: " << fmha_args.p_undrop << "\n";
+    *log_file << "dropout_seed_ptr: " << std::get<0>(std::get<std::pair<const void*, const void*>>(fmha_args.drop_seed_offset)) << "\n";
+    *log_file << "dropout_offset_ptr: " << std::get<1>(std::get<std::pair<const void*, const void*>>(fmha_args.drop_seed_offset)) << "\n";
   }
 
 }
@@ -529,15 +551,10 @@ hipError_t ck_attn_bwd(
   right = window_size_right;
  
   mask_enum mask_type = static_cast<mask_enum>(attn_mask_type);
-  bool ck_fused_attn_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_fused_attn_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, ck_fused_attn_log_config};
+  ck_tile::stream_config stream_config{stream, dump_path!=nullptr, get_bwd_log_stream() != nullptr};
 
   ck_tile::index_t shape_seqlen_q = seqlen_q;
   ck_tile::index_t shape_seqlen_k = seqlen_k;
@@ -707,18 +724,18 @@ hipError_t ck_attn_bwd(
     dim3 grid(b, s_kv, hg);
     if (d_qk == d_v) {
       dim3 block(d_qk);
-      if (ck_fused_attn_log_config){
-        std::cout<<std::endl<<"run dk_dv_reduce: "<<std::endl;
-        std::cout<<"dk_expanded_ptr: "<<dk_expanded_ptr<<std::endl;
-        std::cout<<"dv_expanded_ptr: "<<dv_expanded_ptr<<std::endl;
-        std::cout<<"stride_b_dkv_expanded: "<<stride_b_dk_expanded<<std::endl;
-        std::cout<<"stride_h_dkv_expanded: "<<stride_h_dk_expanded<<std::endl;
-        std::cout<<"stride_s_dkv_expanded: "<<stride_s_dk_expanded<<std::endl;
-        std::cout<<"dk_ptr: "<<dk_ptr<<std::endl;
-        std::cout<<"dv_ptr: "<<dv_ptr<<std::endl;
-        std::cout<<"stride_b_dk: "<<stride_b_dk<<std::endl;
-        std::cout<<"stride_h_dk: "<<stride_h_dk<<std::endl;
-        std::cout<<"stride_s_dk: "<<stride_s_dk<<std::endl;
+      if (auto* log_file = get_bwd_log_stream()) {
+        *log_file << "\n" << "run dk_dv_reduce: " << "\n";
+        *log_file << "dk_expanded_ptr: " << dk_expanded_ptr << "\n";
+        *log_file << "dv_expanded_ptr: " << dv_expanded_ptr << "\n";
+        *log_file << "stride_b_dkv_expanded: " << stride_b_dk_expanded << "\n";
+        *log_file << "stride_h_dkv_expanded: " << stride_h_dk_expanded << "\n";
+        *log_file << "stride_s_dkv_expanded: " << stride_s_dk_expanded << "\n";
+        *log_file << "dk_ptr: " << dk_ptr << "\n";
+        *log_file << "dv_ptr: " << dv_ptr << "\n";
+        *log_file << "stride_b_dk: " << stride_b_dk << "\n";
+        *log_file << "stride_h_dk: " << stride_h_dk << "\n";
+        *log_file << "stride_s_dk: " << stride_s_dk << "\n";
       }
       CK_FUSED_ATTN_TYPE_SWITCH_16BIT(dtype, CK_TILE_TYPE,
         hipLaunchKernelGGL(
@@ -732,16 +749,16 @@ hipError_t ck_attn_bwd(
           stride_b_dk, stride_h_dk, stride_s_dk););
     } else {
       dim3 block_dk(d_qk);
-      if (ck_fused_attn_log_config){
-        std::cout<<std::endl<<"run dk_or_dv_reduce on dk: "<<std::endl;
-        std::cout<<"dk_expanded_ptr: "<<dk_expanded_ptr<<std::endl;
-        std::cout<<"stride_b_dk_expanded: "<<stride_b_dk_expanded<<std::endl;
-        std::cout<<"stride_h_dk_expanded: "<<stride_h_dk_expanded<<std::endl;
-        std::cout<<"stride_s_dk_expanded: "<<stride_s_dk_expanded<<std::endl;
-        std::cout<<"dk_ptr: "<<dk_ptr<<std::endl;
-        std::cout<<"stride_b_dk: "<<stride_b_dk<<std::endl;
-        std::cout<<"stride_h_dk: "<<stride_h_dk<<std::endl;
-        std::cout<<"stride_s_dk: "<<stride_s_dk<<std::endl;
+      if (auto* log_file = get_bwd_log_stream()) {
+        *log_file << "\n" << "run dk_or_dv_reduce on dk: " << "\n";
+        *log_file << "dk_expanded_ptr: " << dk_expanded_ptr << "\n";
+        *log_file << "stride_b_dk_expanded: " << stride_b_dk_expanded << "\n";
+        *log_file << "stride_h_dk_expanded: " << stride_h_dk_expanded << "\n";
+        *log_file << "stride_s_dk_expanded: " << stride_s_dk_expanded << "\n";
+        *log_file << "dk_ptr: " << dk_ptr << "\n";
+        *log_file << "stride_b_dk: " << stride_b_dk << "\n";
+        *log_file << "stride_h_dk: " << stride_h_dk << "\n";
+        *log_file << "stride_s_dk: " << stride_s_dk << "\n";
       }
       CK_FUSED_ATTN_TYPE_SWITCH_16BIT(dtype, CK_TILE_TYPE,
         hipLaunchKernelGGL(
@@ -753,16 +770,16 @@ hipError_t ck_attn_bwd(
           stride_b_dk, stride_h_dk, stride_s_dk););
 
       dim3 block_dv(d_v);
-      if (ck_fused_attn_log_config){
-        std::cout<<std::endl<<"run dk_or_dv_reduce on dv: "<<std::endl;
-        std::cout<<"dv_expanded_ptr: "<<dv_expanded_ptr<<std::endl;
-        std::cout<<"stride_b_dv_expanded: "<<stride_b_dv_expanded<<std::endl;
-        std::cout<<"stride_h_dv_expanded: "<<stride_h_dv_expanded<<std::endl;
-        std::cout<<"stride_s_dv_expanded: "<<stride_s_dv_expanded<<std::endl;
-        std::cout<<"dv_ptr: "<<dv_ptr<<std::endl;
-        std::cout<<"stride_b_dv: "<<stride_b_dv<<std::endl;
-        std::cout<<"stride_h_dv: "<<stride_h_dv<<std::endl;
-        std::cout<<"stride_s_dv: "<<stride_s_dv<<std::endl;
+      if (auto* log_file = get_bwd_log_stream()) {
+        *log_file << "\n" << "run dk_or_dv_reduce on dv: " << "\n";
+        *log_file << "dv_expanded_ptr: " << dv_expanded_ptr << "\n";
+        *log_file << "stride_b_dv_expanded: " << stride_b_dv_expanded << "\n";
+        *log_file << "stride_h_dv_expanded: " << stride_h_dv_expanded << "\n";
+        *log_file << "stride_s_dv_expanded: " << stride_s_dv_expanded << "\n";
+        *log_file << "dv_ptr: " << dv_ptr << "\n";
+        *log_file << "stride_b_dv: " << stride_b_dv << "\n";
+        *log_file << "stride_h_dv: " << stride_h_dv << "\n";
+        *log_file << "stride_s_dv: " << stride_s_dv << "\n";
       }
       CK_FUSED_ATTN_TYPE_SWITCH_16BIT(dtype, CK_TILE_TYPE,
         hipLaunchKernelGGL(
@@ -781,10 +798,10 @@ hipError_t ck_attn_bwd(
     dim3 block(THREADS_PER_BLOCK);
     dim3 grid(ceil(1.0 * s_q * s_kv/THREADS_PER_BLOCK));
     if(bias_shape==BiasShape::k11SS){
-      if (ck_fused_attn_log_config){
-        std::cout<<std::endl<<"run dbias_reduce_11SS: "<<std::endl;
-        std::cout<<"dbias_ptr: "<<dbias_ptr<<std::endl;
-        std::cout<<"dbias_expanded_ptr: "<<dbias_expanded_ptr<<std::endl;
+      if (auto* log_file = get_bwd_log_stream()) {
+        *log_file << "\n" << "run dbias_reduce_11SS: " << "\n";
+        *log_file << "dbias_ptr: " << dbias_ptr << "\n";
+        *log_file << "dbias_expanded_ptr: " << dbias_expanded_ptr << "\n";
       }
       CK_FUSED_ATTN_TYPE_SWITCH_16BIT(dtype, CK_TILE_TYPE,
         hipLaunchKernelGGL(
@@ -793,10 +810,10 @@ hipError_t ck_attn_bwd(
           static_cast<CK_TILE_TYPE*>(dbias_expanded_ptr),
           static_cast<CK_TILE_TYPE*>(dbias_ptr));); 
     }else if(bias_shape==BiasShape::k1HSS){
-      if (ck_fused_attn_log_config){
-        std::cout<<std::endl<<"run dbias_reduce_1HSS: "<<std::endl;
-        std::cout<<"dbias_ptr: "<<dbias_ptr<<std::endl;
-        std::cout<<"dbias_expanded_ptr: "<<dbias_expanded_ptr<<std::endl;
+      if (auto* log_file = get_bwd_log_stream()) {
+        *log_file << "\n" << "run dbias_reduce_1HSS: " << "\n";
+        *log_file << "dbias_ptr: " << dbias_ptr << "\n";
+        *log_file << "dbias_expanded_ptr: " << dbias_expanded_ptr << "\n";
       }
       CK_FUSED_ATTN_TYPE_SWITCH_16BIT(dtype, CK_TILE_TYPE,
         hipLaunchKernelGGL(
@@ -805,10 +822,10 @@ hipError_t ck_attn_bwd(
           static_cast<CK_TILE_TYPE*>(dbias_expanded_ptr),
           static_cast<CK_TILE_TYPE*>(dbias_ptr));); 
     }else if(bias_shape==BiasShape::kB1SS){
-      if (ck_fused_attn_log_config){
-        std::cout<<std::endl<<"run dbias_reduce_B1SS: "<<std::endl;
-        std::cout<<"dbias_ptr: "<<dbias_ptr<<std::endl;
-        std::cout<<"dbias_expanded_ptr: "<<dbias_expanded_ptr<<std::endl;
+      if (auto* log_file = get_bwd_log_stream()) {
+        *log_file << "\n" << "run dbias_reduce_B1SS: " << "\n";
+        *log_file << "dbias_ptr: " << dbias_ptr << "\n";
+        *log_file << "dbias_expanded_ptr: " << dbias_expanded_ptr << "\n";
       }
       CK_FUSED_ATTN_TYPE_SWITCH_16BIT(dtype, CK_TILE_TYPE,
         hipLaunchKernelGGL(
@@ -885,14 +902,9 @@ hipError_t ck_attn_varlen_bwd(
   right = window_size_right;
   mask_enum mask_type = static_cast<mask_enum>(attn_mask_type);
  
-  bool ck_fused_attn_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_fused_attn_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, ck_fused_attn_log_config};
+  ck_tile::stream_config stream_config{stream, dump_path!=nullptr, get_bwd_log_stream() != nullptr};
 
   std::string data_type_str = get_data_type_str(dtype);
 
@@ -1034,8 +1046,9 @@ hipError_t ck_attn_varlen_bwd(
   // lse_thd_ptr used as buffer
   if(const char* env_p = std::getenv("NVTE_CK_RUNTIME_MAX_SEQLEN")) {
     if(std::string(env_p) == "1"){
-      if(ck_fused_attn_log_config){
-        std::cout << "attn_bwd(ck): Enabling runtime max_seqlen calculation for small seqlen optimization.";
+      if (auto* log_file = get_bwd_log_stream()) {
+        *log_file
+            << "attn_bwd(ck): Enabling runtime max_seqlen calculation for small seqlen optimization.\n";
       }
       fmha_args.max_seqlen_q = get_runtime_max_seqlen(b, cu_seqlen_q_ptr, nullptr, lse_workspace_ptr, stream);
       fmha_args.max_seqlen_k = get_runtime_max_seqlen(b, cu_seqlen_kv_ptr, nullptr, lse_workspace_ptr, stream);
@@ -1068,18 +1081,18 @@ hipError_t ck_attn_varlen_bwd(
     dim3 grid(max_tokens_kv, hg);
     if (d_qk == d_v) {
       dim3 block(d_qk);
-      if (ck_fused_attn_log_config){
-        std::cout<<std::endl<<"run dk_dv_reduce_thd: "<<std::endl;
-        std::cout<<"cu_seqlen_kv_ptr: "<<cu_seqlen_kv_ptr<<std::endl;
-        std::cout<<"cu_seqlen_kv_padded_ptr: "<<cu_seqlen_kv_padded_ptr<<std::endl;
-        std::cout<<"dk_expanded_ptr: "<<dk_expanded_ptr<<std::endl;
-        std::cout<<"dv_expanded_ptr: "<<dv_expanded_ptr<<std::endl;
-        std::cout<<"stride_h_dkv_expanded: "<<stride_h_dk_expanded<<std::endl;
-        std::cout<<"stride_s_dkv_expanded: "<<stride_s_dk_expanded<<std::endl;
-        std::cout<<"dk_ptr: "<<dk_ptr<<std::endl;
-        std::cout<<"dv_ptr: "<<dv_ptr<<std::endl;
-        std::cout<<"stride_h_dk: "<<stride_h_dk<<std::endl;
-        std::cout<<"stride_s_dk: "<<stride_s_dk<<std::endl;
+      if (auto* log_file = get_bwd_log_stream()) {
+        *log_file << "\n" << "run dk_dv_reduce_thd: " << "\n";
+        *log_file << "cu_seqlen_kv_ptr: " << cu_seqlen_kv_ptr << "\n";
+        *log_file << "cu_seqlen_kv_padded_ptr: " << cu_seqlen_kv_padded_ptr << "\n";
+        *log_file << "dk_expanded_ptr: " << dk_expanded_ptr << "\n";
+        *log_file << "dv_expanded_ptr: " << dv_expanded_ptr << "\n";
+        *log_file << "stride_h_dkv_expanded: " << stride_h_dk_expanded << "\n";
+        *log_file << "stride_s_dkv_expanded: " << stride_s_dk_expanded << "\n";
+        *log_file << "dk_ptr: " << dk_ptr << "\n";
+        *log_file << "dv_ptr: " << dv_ptr << "\n";
+        *log_file << "stride_h_dk: " << stride_h_dk << "\n";
+        *log_file << "stride_s_dk: " << stride_s_dk << "\n";
       }
       CK_FUSED_ATTN_TYPE_SWITCH_16BIT(dtype, CK_TILE_TYPE,
         hipLaunchKernelGGL(
@@ -1095,16 +1108,16 @@ hipError_t ck_attn_varlen_bwd(
           stride_h_dk, stride_s_dk););
     } else {
       dim3 block_dk(d_qk);
-      if (ck_fused_attn_log_config){
-        std::cout<<std::endl<<"run dk_or_dv_reduce_thd on dk: "<<std::endl;
-        std::cout<<"cu_seqlen_kv_ptr: "<<cu_seqlen_kv_ptr<<std::endl;
-        std::cout<<"cu_seqlen_kv_padded_ptr: "<<cu_seqlen_kv_padded_ptr<<std::endl;
-        std::cout<<"dk_expanded_ptr: "<<dk_expanded_ptr<<std::endl;
-        std::cout<<"stride_h_dk_expanded: "<<stride_h_dk_expanded<<std::endl;
-        std::cout<<"stride_s_dk_expanded: "<<stride_s_dk_expanded<<std::endl;
-        std::cout<<"dk_ptr: "<<dk_ptr<<std::endl;
-        std::cout<<"stride_h_dk: "<<stride_h_dk<<std::endl;
-        std::cout<<"stride_s_dk: "<<stride_s_dk<<std::endl;
+      if (auto* log_file = get_bwd_log_stream()) {
+        *log_file << "\n" << "run dk_or_dv_reduce_thd on dk: " << "\n";
+        *log_file << "cu_seqlen_kv_ptr: " << cu_seqlen_kv_ptr << "\n";
+        *log_file << "cu_seqlen_kv_padded_ptr: " << cu_seqlen_kv_padded_ptr << "\n";
+        *log_file << "dk_expanded_ptr: " << dk_expanded_ptr << "\n";
+        *log_file << "stride_h_dk_expanded: " << stride_h_dk_expanded << "\n";
+        *log_file << "stride_s_dk_expanded: " << stride_s_dk_expanded << "\n";
+        *log_file << "dk_ptr: " << dk_ptr << "\n";
+        *log_file << "stride_h_dk: " << stride_h_dk << "\n";
+        *log_file << "stride_s_dk: " << stride_s_dk << "\n";
       }
       CK_FUSED_ATTN_TYPE_SWITCH_16BIT(dtype, CK_TILE_TYPE,
         hipLaunchKernelGGL(
@@ -1118,16 +1131,16 @@ hipError_t ck_attn_varlen_bwd(
           stride_h_dk, stride_s_dk););
 
       dim3 block_dv(d_v);
-      if (ck_fused_attn_log_config){
-        std::cout<<std::endl<<"run dk_or_dv_reduce_thd on dv: "<<std::endl;
-        std::cout<<"cu_seqlen_kv_ptr: "<<cu_seqlen_kv_ptr<<std::endl;
-        std::cout<<"cu_seqlen_kv_padded_ptr: "<<cu_seqlen_kv_padded_ptr<<std::endl;
-        std::cout<<"dv_expanded_ptr: "<<dv_expanded_ptr<<std::endl;
-        std::cout<<"stride_h_dv_expanded: "<<stride_h_dv_expanded<<std::endl;
-        std::cout<<"stride_s_dv_expanded: "<<stride_s_dv_expanded<<std::endl;
-        std::cout<<"dv_ptr: "<<dv_ptr<<std::endl;
-        std::cout<<"stride_h_dv: "<<stride_h_dv<<std::endl;
-        std::cout<<"stride_s_dv: "<<stride_s_dv<<std::endl;
+      if (auto* log_file = get_bwd_log_stream()) {
+        *log_file << "\n" << "run dk_or_dv_reduce_thd on dv: " << "\n";
+        *log_file << "cu_seqlen_kv_ptr: " << cu_seqlen_kv_ptr << "\n";
+        *log_file << "cu_seqlen_kv_padded_ptr: " << cu_seqlen_kv_padded_ptr << "\n";
+        *log_file << "dv_expanded_ptr: " << dv_expanded_ptr << "\n";
+        *log_file << "stride_h_dv_expanded: " << stride_h_dv_expanded << "\n";
+        *log_file << "stride_s_dv_expanded: " << stride_s_dv_expanded << "\n";
+        *log_file << "dv_ptr: " << dv_ptr << "\n";
+        *log_file << "stride_h_dv: " << stride_h_dv << "\n";
+        *log_file << "stride_s_dv: " << stride_s_dv << "\n";
       }
       CK_FUSED_ATTN_TYPE_SWITCH_16BIT(dtype, CK_TILE_TYPE,
         hipLaunchKernelGGL(
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 8d7c4b5c9..935caa9eb 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
@@ -1,5 +1,5 @@
 /*************************************************************************
- * Copyright (c) 2024-2025, Advanced Micro Devices, Inc. All rights reserved.
+ * Copyright (c) 2024-2026, Advanced Micro Devices, Inc. All rights reserved.
  *
  * License for AMD contributions = MIT. See LICENSE for more information
  ************************************************************************/
@@ -16,6 +16,33 @@
 
 namespace ck_fused_attn{
 
+namespace {
+std::ostream* get_fwd_log_stream() {
+  thread_local std::ofstream log_file;
+  thread_local bool attempted = false;
+  thread_local bool opened = false;
+  thread_local bool requested = false;
+  thread_local std::string log_dir_str;
+  if (!attempted) {
+    attempted = true;
+    if (const char* env_p = std::getenv("CK_FUSED_ATTN_LOG_CONFIG")) {
+      log_dir_str = std::string(env_p);
+      requested = !log_dir_str.empty() && log_dir_str != "0";
+    }
+    if (requested) {
+      opened = open_ck_fused_attn_log_file(log_file, "ck_fused_attn_fwd", log_dir_str);
+    }
+  }
+  if (!requested) {
+    return nullptr;
+  }
+  if (!opened) {
+    return &std::cout;
+  }
+  return &log_file;
+}
+}  // namespace
+
 // print the fmha traits and args when calling ck apis
 void log_fwd_config(const char* func_name,
                     const std::string data_type_str,
@@ -30,93 +57,88 @@ void log_fwd_config(const char* func_name,
                     const bool uses_fwd_v3,
                     const bool how_v3_bf16_cvt,
                     const fmha_fwd_args& fmha_args){
-  bool ck_fused_attn_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_fused_attn_log_config = true;
-  }
-  if (ck_fused_attn_log_config) {
-    std::cout<<std::endl<<func_name<<std::endl;
+  if (auto* log_file = get_fwd_log_stream()) {
+    *log_file << "\n" << func_name << "\n";
 
     // debug fmha_traits
-    std::cout<<std::endl<<"fmha_traits: "<<std::endl;
-    std::cout<<"hdim_q: "<<fmha_args.hdim_q<<std::endl;
-    std::cout<<"hdim_v: "<<fmha_args.hdim_v<<std::endl;
-    std::cout<<"data_type: "<<data_type_str<<std::endl;
-    std::cout<<"is_group_mode: "<<is_group_mode<<std::endl;
-    std::cout<<"is_v_rowmajor: "<<is_v_rowmajor<<std::endl;
-    std::cout<<"has_logits_soft_cap: "<<has_logits_soft_cap<<std::endl;
-    std::cout<<"mask_type: "<<static_cast<std::underlying_type<mask_enum>::type>(mask_type)<<std::endl;
-    std::cout<<"bias_type: "<<static_cast<std::underlying_type<bias_enum>::type>(bias_type)<<std::endl;
-    std::cout<<"has_lse: "<<has_lse<<std::endl;
-    std::cout<<"has_dropout: "<<has_dropout<<std::endl;
-    std::cout<<"do_fp8_static_quant: "<<do_fp8_static_quant<<std::endl;
-    std::cout<<"skip_min_seqlen_q: "<<(fmha_args.min_seqlen_q != 0)<<std::endl;
-    std::cout<<"uses_fwd_v3: "<<uses_fwd_v3<<std::endl;
-    std::cout<<"how_v3_bf16_cvt: "<<how_v3_bf16_cvt<<std::endl;
+    *log_file << "\n" << "fmha_traits: " << "\n";
+    *log_file << "hdim_q: " << fmha_args.hdim_q << "\n";
+    *log_file << "hdim_v: " << fmha_args.hdim_v << "\n";
+    *log_file << "data_type: " << data_type_str << "\n";
+    *log_file << "is_group_mode: " << is_group_mode << "\n";
+    *log_file << "is_v_rowmajor: " << is_v_rowmajor << "\n";
+    *log_file << "has_logits_soft_cap: " << has_logits_soft_cap << "\n";
+    *log_file << "mask_type: " << static_cast<std::underlying_type<mask_enum>::type>(mask_type) << "\n";
+    *log_file << "bias_type: " << static_cast<std::underlying_type<bias_enum>::type>(bias_type) << "\n";
+    *log_file << "has_lse: " << has_lse << "\n";
+    *log_file << "has_dropout: " << has_dropout << "\n";
+    *log_file << "do_fp8_static_quant: " << do_fp8_static_quant << "\n";
+    *log_file << "skip_min_seqlen_q: " << (fmha_args.min_seqlen_q != 0) << "\n";
+    *log_file << "uses_fwd_v3: " << uses_fwd_v3 << "\n";
+    *log_file << "how_v3_bf16_cvt: " << how_v3_bf16_cvt << "\n";
 
     // debug fmha_args
-    std::cout<<std::endl<<"fmha_args: "<<std::endl;
-
-    std::cout<<"q_ptr: "<<fmha_args.q_ptr<<std::endl;
-    std::cout<<"k_ptr: "<<fmha_args.k_ptr<<std::endl;
-    std::cout<<"v_ptr: "<<fmha_args.v_ptr<<std::endl;
-    std::cout<<"bias_ptr: "<<fmha_args.bias_ptr<<std::endl;
-    std::cout<<"rand_val_ptr: "<<fmha_args.rand_val_ptr<<std::endl;
-    std::cout<<"lse_ptr: "<<fmha_args.lse_ptr<<std::endl;
-    std::cout<<"o_ptr: "<<fmha_args.o_ptr<<std::endl;
-
-    std::cout<<"seqstart_q_ptr: "<<fmha_args.seqstart_q_ptr<<std::endl;
-    std::cout<<"seqstart_k_ptr: "<<fmha_args.seqstart_k_ptr<<std::endl;
-    std::cout<<"seqlen_q_ptr: "<<fmha_args.seqlen_q_ptr<<std::endl;
-    std::cout<<"seqlen_k_ptr: "<<fmha_args.seqlen_k_ptr<<std::endl;
-    std::cout<<"cu_seqlen_q_ptr: "<<fmha_args.cu_seqlen_q_ptr<<std::endl;
-    std::cout<<"cu_seqlen_k_ptr: "<<fmha_args.cu_seqlen_k_ptr<<std::endl;
-
-    std::cout<<"seqlen_q: "<<fmha_args.seqlen_q<<std::endl;
-    std::cout<<"seqlen_k: "<<fmha_args.seqlen_k<<std::endl;
-    std::cout<<"batch: "<<fmha_args.batch<<std::endl;
-    std::cout<<"max_seqlen_q: "<<fmha_args.max_seqlen_q<<std::endl;
-    std::cout<<"hdim_q: "<<fmha_args.hdim_q<<std::endl;
-    std::cout<<"hdim_v: "<<fmha_args.hdim_v<<std::endl;
-    std::cout<<"nhead_q: "<<fmha_args.nhead_q<<std::endl;
-    std::cout<<"nhead_k: "<<fmha_args.nhead_k<<std::endl;
-
-    std::cout<<"scale_s: "<<fmha_args.scale_s<<std::endl;
-
-    std::cout<<"logits_soft_cap: "<<fmha_args.logits_soft_cap<<std::endl;
-
-    std::cout<<"stride_q: "<<fmha_args.stride_q<<std::endl;
-    std::cout<<"stride_k: "<<fmha_args.stride_k<<std::endl;
-    std::cout<<"stride_v: "<<fmha_args.stride_v<<std::endl;
-    std::cout<<"stride_bias: "<<fmha_args.stride_bias<<std::endl;
-    std::cout<<"stride_randval: "<<fmha_args.stride_randval<<std::endl;
-    std::cout<<"stride_o: "<<fmha_args.stride_o<<std::endl;
-    std::cout<<"nhead_stride_q: "<<fmha_args.nhead_stride_q<<std::endl;
-    std::cout<<"nhead_stride_k: "<<fmha_args.nhead_stride_k<<std::endl;
-    std::cout<<"nhead_stride_v: "<<fmha_args.nhead_stride_v<<std::endl;
-    std::cout<<"nhead_stride_bias: "<<fmha_args.nhead_stride_bias<<std::endl;
-    std::cout<<"nhead_stride_randval: "<<fmha_args.nhead_stride_randval<<std::endl;
-    std::cout<<"nhead_stride_lse: "<<fmha_args.nhead_stride_lse<<std::endl;
-    std::cout<<"nhead_stride_o: "<<fmha_args.nhead_stride_o<<std::endl;
-    std::cout<<"batch_stride_q: "<<fmha_args.batch_stride_q<<std::endl;
-    std::cout<<"batch_stride_k: "<<fmha_args.batch_stride_k<<std::endl;
-    std::cout<<"batch_stride_v: "<<fmha_args.batch_stride_v<<std::endl;
-    std::cout<<"batch_stride_bias: "<<fmha_args.batch_stride_bias<<std::endl;
-    std::cout<<"batch_stride_randval: "<<fmha_args.batch_stride_randval<<std::endl;
-    std::cout<<"batch_stride_lse: "<<fmha_args.batch_stride_lse<<std::endl;
-    std::cout<<"batch_stride_o: "<<fmha_args.batch_stride_o<<std::endl;
-
-    std::cout<<"window_size_left: "<<fmha_args.window_size_left<<std::endl;
-    std::cout<<"window_size_right: "<<fmha_args.window_size_right<<std::endl;
-    std::cout<<"mask_type: "<<fmha_args.mask_type<<std::endl;
-    std::cout<<"min_seqlen_q: "<<fmha_args.min_seqlen_q<<std::endl;
-
-    std::cout<<"p_drop: "<<fmha_args.p_drop<<std::endl;
-    std::cout<<"s_randval: "<<fmha_args.s_randval<<std::endl;
-
-    std::cout<<"dropout_seed_ptr: "<<std::get<0>(std::get<std::pair<const void*, const void*>>(fmha_args.drop_seed_offset))<<std::endl;
-    std::cout<<"dropout_offset_ptr: "<<std::get<1>(std::get<std::pair<const void*, const void*>>(fmha_args.drop_seed_offset))<<std::endl;
+    *log_file << "\n" << "fmha_args: " << "\n";
+
+    *log_file << "q_ptr: " << fmha_args.q_ptr << "\n";
+    *log_file << "k_ptr: " << fmha_args.k_ptr << "\n";
+    *log_file << "v_ptr: " << fmha_args.v_ptr << "\n";
+    *log_file << "bias_ptr: " << fmha_args.bias_ptr << "\n";
+    *log_file << "rand_val_ptr: " << fmha_args.rand_val_ptr << "\n";
+    *log_file << "lse_ptr: " << fmha_args.lse_ptr << "\n";
+    *log_file << "o_ptr: " << fmha_args.o_ptr << "\n";
+
+    *log_file << "seqstart_q_ptr: " << fmha_args.seqstart_q_ptr << "\n";
+    *log_file << "seqstart_k_ptr: " << fmha_args.seqstart_k_ptr << "\n";
+    *log_file << "seqlen_q_ptr: " << fmha_args.seqlen_q_ptr << "\n";
+    *log_file << "seqlen_k_ptr: " << fmha_args.seqlen_k_ptr << "\n";
+    *log_file << "cu_seqlen_q_ptr: " << fmha_args.cu_seqlen_q_ptr << "\n";
+    *log_file << "cu_seqlen_k_ptr: " << fmha_args.cu_seqlen_k_ptr << "\n";
+
+    *log_file << "seqlen_q: " << fmha_args.seqlen_q << "\n";
+    *log_file << "seqlen_k: " << fmha_args.seqlen_k << "\n";
+    *log_file << "batch: " << fmha_args.batch << "\n";
+    *log_file << "max_seqlen_q: " << fmha_args.max_seqlen_q << "\n";
+    *log_file << "hdim_q: " << fmha_args.hdim_q << "\n";
+    *log_file << "hdim_v: " << fmha_args.hdim_v << "\n";
+    *log_file << "nhead_q: " << fmha_args.nhead_q << "\n";
+    *log_file << "nhead_k: " << fmha_args.nhead_k << "\n";
+
+    *log_file << "scale_s: " << fmha_args.scale_s << "\n";
+
+    *log_file << "logits_soft_cap: " << fmha_args.logits_soft_cap << "\n";
+
+    *log_file << "stride_q: " << fmha_args.stride_q << "\n";
+    *log_file << "stride_k: " << fmha_args.stride_k << "\n";
+    *log_file << "stride_v: " << fmha_args.stride_v << "\n";
+    *log_file << "stride_bias: " << fmha_args.stride_bias << "\n";
+    *log_file << "stride_randval: " << fmha_args.stride_randval << "\n";
+    *log_file << "stride_o: " << fmha_args.stride_o << "\n";
+    *log_file << "nhead_stride_q: " << fmha_args.nhead_stride_q << "\n";
+    *log_file << "nhead_stride_k: " << fmha_args.nhead_stride_k << "\n";
+    *log_file << "nhead_stride_v: " << fmha_args.nhead_stride_v << "\n";
+    *log_file << "nhead_stride_bias: " << fmha_args.nhead_stride_bias << "\n";
+    *log_file << "nhead_stride_randval: " << fmha_args.nhead_stride_randval << "\n";
+    *log_file << "nhead_stride_lse: " << fmha_args.nhead_stride_lse << "\n";
+    *log_file << "nhead_stride_o: " << fmha_args.nhead_stride_o << "\n";
+    *log_file << "batch_stride_q: " << fmha_args.batch_stride_q << "\n";
+    *log_file << "batch_stride_k: " << fmha_args.batch_stride_k << "\n";
+    *log_file << "batch_stride_v: " << fmha_args.batch_stride_v << "\n";
+    *log_file << "batch_stride_bias: " << fmha_args.batch_stride_bias << "\n";
+    *log_file << "batch_stride_randval: " << fmha_args.batch_stride_randval << "\n";
+    *log_file << "batch_stride_lse: " << fmha_args.batch_stride_lse << "\n";
+    *log_file << "batch_stride_o: " << fmha_args.batch_stride_o << "\n";
+
+    *log_file << "window_size_left: " << fmha_args.window_size_left << "\n";
+    *log_file << "window_size_right: " << fmha_args.window_size_right << "\n";
+    *log_file << "mask_type: " << fmha_args.mask_type << "\n";
+    *log_file << "min_seqlen_q: " << fmha_args.min_seqlen_q << "\n";
+
+    *log_file << "p_drop: " << fmha_args.p_drop << "\n";
+    *log_file << "s_randval: " << fmha_args.s_randval << "\n";
+
+    *log_file << "dropout_seed_ptr: " << std::get<0>(std::get<std::pair<const void*, const void*>>(fmha_args.drop_seed_offset)) << "\n";
+    *log_file << "dropout_offset_ptr: " << std::get<1>(std::get<std::pair<const void*, const void*>>(fmha_args.drop_seed_offset)) << "\n";
   }
 }
 
@@ -179,14 +201,9 @@ hipError_t ck_attn_fwd(
   right = window_size_right;
   mask_enum mask_type = static_cast<mask_enum>(attn_mask_type);
   
-  bool ck_fused_attn_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_fused_attn_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, ck_fused_attn_log_config};
+  ck_tile::stream_config stream_config{stream, dump_path!=nullptr, get_fwd_log_stream() != nullptr};
 
   std::string data_type_str = get_data_type_str(dtype);
 
@@ -354,14 +371,9 @@ hipError_t ck_attn_varlen_fwd(
   
   bias_enum bias_type = bias_enum::no_bias;
   
-  bool ck_fused_attn_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_fused_attn_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, ck_fused_attn_log_config};
+  ck_tile::stream_config stream_config{stream, dump_path!=nullptr, get_fwd_log_stream() != nullptr};
 
 
   std::string data_type_str = get_data_type_str(dtype);
@@ -457,8 +469,9 @@ hipError_t ck_attn_varlen_fwd(
   // lse_thd_ptr used as buffer
   if(const char* env_p = std::getenv("NVTE_CK_RUNTIME_MAX_SEQLEN")){
     if(std::string(env_p) == "1"){
-      if(ck_fused_attn_log_config){
-        std::cout << "attn_fwd(ck): Enabling runtime max_seqlen calculation for small seqlen optimization.";
+      if (auto* log_file = get_fwd_log_stream()) {
+        *log_file
+            << "attn_fwd(ck): Enabling runtime max_seqlen calculation for small seqlen optimization.\n";
       }
       fmha_args.max_seqlen_q = get_runtime_max_seqlen(b, cu_seqlen_q_ptr, cu_seqlen_q_padded_ptr, lse_thd_ptr, stream);
     }
diff --git a/transformer_engine/common/ck_fused_attn/src/ck_fused_attn_utils.cpp b/transformer_engine/common/ck_fused_attn/src/ck_fused_attn_utils.cpp
index 26c92ca2b..6bbfbda4f 100644
--- a/transformer_engine/common/ck_fused_attn/src/ck_fused_attn_utils.cpp
+++ b/transformer_engine/common/ck_fused_attn/src/ck_fused_attn_utils.cpp
@@ -1,10 +1,14 @@
 /*************************************************************************
- * Copyright (c) 2024-2025, Advanced Micro Devices, Inc. All rights reserved.
+ * Copyright (c) 2024-2026, Advanced Micro Devices, Inc. All rights reserved.
  *
  * License for AMD contributions = MIT. See LICENSE for more information
  ************************************************************************/
 
 #include <utility>
+#include <filesystem>
+#include <sstream>
+#include <thread>
+#include <unistd.h>
 #include "ck_fused_attn_utils.hpp"
 #include "ck_fused_attn/ck_fused_attn.hpp"
 #include "mask.hpp"
@@ -13,6 +17,22 @@
 
 namespace ck_fused_attn{
 
+bool open_ck_fused_attn_log_file(std::ofstream& log_file, const char* file_prefix, const std::string& log_dir_str) {
+  // Explicitly use std::cout as a fallback
+  if (log_dir_str == "1") {
+    return false;
+  }
+  std::filesystem::path log_dir(log_dir_str);
+  std::ostringstream filename;
+  filename << file_prefix << "_" << getpid() << "_" << std::this_thread::get_id() << ".log";
+  log_file.open(log_dir / filename.str(), std::ios_base::app);
+  if (!log_file.is_open()) {
+    std::cerr << "Failed to open log file: " << (log_dir / filename.str()) << "\n";
+    return false;
+  }
+  return true;
+}
+
 std::string get_data_type_str(DType dtype){
   std::string data_type_str;
   if(dtype==DType::kFloat16){
diff --git a/transformer_engine/common/ck_fused_attn/src/ck_fused_attn_utils.hpp b/transformer_engine/common/ck_fused_attn/src/ck_fused_attn_utils.hpp
index a75915ee2..a0ea13d81 100644
--- a/transformer_engine/common/ck_fused_attn/src/ck_fused_attn_utils.hpp
+++ b/transformer_engine/common/ck_fused_attn/src/ck_fused_attn_utils.hpp
@@ -1,5 +1,5 @@
 /*************************************************************************
- * Copyright (c) 2024-2025, Advanced Micro Devices, Inc. All rights reserved.
+ * Copyright (c) 2024-2026, Advanced Micro Devices, Inc. All rights reserved.
  *
  * License for AMD contributions = MIT. See LICENSE for more information
  ************************************************************************/
@@ -7,8 +7,9 @@
 #ifndef CK_FUSED_ATTN_UTILS_H
 #define CK_FUSED_ATTN_UTILS_H
 
-#include<iostream>
-#include<cstdint>
+#include<cstdlib>
+#include<fstream>
+#include<string>
 #include<hip/hip_runtime.h>
 
 //forward declaration for ck_tile enum
@@ -56,5 +57,7 @@ std::pair<bias_enum, BiasShape> get_ck_bias_type_shape(BiasType attn_bias_type,
 
 uint64_t get_runtime_max_seqlen(uint64_t b, const void* cu_seqlen_ptr, const void* cu_seqlen_padded_ptr, void* workspace, hipStream_t stream);
 
+bool open_ck_fused_attn_log_file(std::ofstream& log_file, const char* file_prefix, const std::string& log_dir_str);
+
 }//namespace ck_fused_attn
 #endif // CK_FUSED_ATTN_UTILS_H