diff --git a/include/ck_tile/core/arch/amd_buffer_coherence.hpp b/include/ck_tile/core/arch/amd_buffer_coherence.hpp
index 2e38c77291..27d2799a59 100644
--- a/include/ck_tile/core/arch/amd_buffer_coherence.hpp
+++ b/include/ck_tile/core/arch/amd_buffer_coherence.hpp
@@ -83,6 +83,9 @@ enum struct amd_buffer_coherence_enum
     slc     = SYSTEM_NT,
     glc_slc = DEVICE_NT | SYSTEM_NT,
 
+    // To forward compatible
+    WAVE_NT0   = INT_MAX,
+    WAVE_NT1   = INT_MAX - 1,
 // gfx94: bit 0 = sc0, bit 1 = nt, bit 3 = swz, bit 4 = sc1
 // SC[1:0] System Cache level: 0=wave, 1=group, 2=device, 3=system
 // NT Non-Temporal: 0=expect temporal reuse; 1=do not expect temporal reuse
@@ -118,6 +121,12 @@ enum struct amd_buffer_coherence_enum
     SYSTEM_NT0 = 0,
     DEVICE_NT1 = glc,
     SYSTEM_NT1 = slc,
+
+    // To forward compatible
+    GROUP_NT0  = INT_MAX - 2,
+    GROUP_NT1  = INT_MAX - 3,
+    WAVE_NT0   = INT_MAX,
+    WAVE_NT1   = INT_MAX - 1,
 #endif
 };
 
diff --git a/include/ck_tile/ops/flatmm/pipeline/flatmm_pipeline_agmem_bgmem_creg_v1.hpp b/include/ck_tile/ops/flatmm/pipeline/flatmm_pipeline_agmem_bgmem_creg_v1.hpp
index 897892afb2..dc80b3964d 100644
--- a/include/ck_tile/ops/flatmm/pipeline/flatmm_pipeline_agmem_bgmem_creg_v1.hpp
+++ b/include/ck_tile/ops/flatmm/pipeline/flatmm_pipeline_agmem_bgmem_creg_v1.hpp
@@ -31,9 +31,7 @@ struct BaseFlatmmPipelineAGmemBGmemCRegV1
         ck_tile::ignore = K;
         if(M <= 416)
         {
-#if defined(__gfx942__) || defined(__gfx950__)
             return ck_tile::amd_buffer_coherence_enum::WAVE_NT1;
-#endif
         }
         return ck_tile::amd_buffer_coherence_enum::coherence_default;
     }