diff --git a/projects/composablekernel/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_blockscale_b_preshuffle_v1.hpp b/projects/composablekernel/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_blockscale_b_preshuffle_v1.hpp
index af9ff1e16aae..214e21385fad 100644
--- a/projects/composablekernel/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_blockscale_b_preshuffle_v1.hpp
+++ b/projects/composablekernel/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_blockscale_b_preshuffle_v1.hpp
@@ -285,6 +285,15 @@ struct BlockwiseGemmXdlops_pipeline_blockscale_bpreshuffle_v1<BlockGemmPipelineS
     {
         ignore = b_block_desc;
         ignore = b_block_buf;
+
+        static_assert(sizeof(AccDataType) == 4 && is_floating_point<AccDataType>::value,
+                      "bpreshuffle accumulator anchor requires 32-bit float AccDataType");
+
+        // The empty asm is a read/write VGPR use of the updated accumulator.
+        // It keeps each post-scale accumulator definition visible to the
+        // optimizer, which enables determinism for bitwise-stable repeated launches.
+        auto anchor_accumulator_value = [&](auto& value) { asm volatile("" : "+v"(value)); };
+
         // __builtin_amdgcn_sched_barrier(0);
         auto a_thread_buf = make_static_buffer<AddressSpaceEnum::Vgpr, ComputeDataType>(
             a_thread_desc_.GetElementSpaceSize());
@@ -451,14 +460,9 @@ struct BlockwiseGemmXdlops_pipeline_blockscale_bpreshuffle_v1<BlockGemmPipelineS
                             c_thread_buf_per_scale.GetVectorTypeReference(Number<0>{})
                                 .template AsType<AccDataType>()(Number<t>{}) = 0;
                         });
-                        vector_type<AccDataType, 2> c_scale_thread_vec;
                         constexpr index_t cscale_offset = CScaleThreadDesc{}.CalculateOffset(
                             make_tuple(kscale0, m0, n0 * num_scale_n_block / NRepeat));
-
-                        c_scale_thread_vec.template AsType<AccDataType>()(Number<0>{}) =
-                            c_scale_thread_buf[Number<cscale_offset>{}];
-                        c_scale_thread_vec.template AsType<AccDataType>()(Number<1>{}) =
-                            c_scale_thread_buf[Number<cscale_offset>{}];
+                        const auto c_scale_thread = c_scale_thread_buf[Number<cscale_offset>{}];
 
                         static_for<0, KRepeat / num_scale_k_block, 1>{}([&](auto k0) {
                             vector_type<ComputeDataType, KPack> a_thread_vec;
@@ -501,16 +505,18 @@ struct BlockwiseGemmXdlops_pipeline_blockscale_bpreshuffle_v1<BlockGemmPipelineS
                         constexpr index_t c_offset =
                             c_thread_desc_.CalculateOffset(make_tuple(m0, n0, 0));
 
-                        static_for<0, xdlops_gemm.GetRegSizePerXdlops() / 2, 1>{}([&](auto t) {
-                            using pk_fma_type = typename vector_type<AccDataType, 2>::type;
-
-                            c_thread_buf.GetVectorTypeReference(Number<c_offset>{})
-                                .template AsType<pk_fma_type>()(t) = __builtin_elementwise_fma(
-                                c_thread_buf_per_scale.GetVectorTypeReference(Number<0>{})
-                                    .template AsType<pk_fma_type>()[t],
-                                c_scale_thread_vec.template AsType<pk_fma_type>()[Number<0>{}],
+                        // Keep the post-scale FMA scalar; the old packed 2-lane update hits
+                        // ROCm 7.2 gfx950 illegal-type legalization/codegen issues.
+                        static_for<0, xdlops_gemm.GetRegSizePerXdlops(), 1>{}([&](auto t) {
+                            auto& c_acc_vec =
                                 c_thread_buf.GetVectorTypeReference(Number<c_offset>{})
-                                    .template AsType<pk_fma_type>()[t]);
+                                    .template AsType<AccDataType>();
+                            const auto c_partial_acc =
+                                c_thread_buf_per_scale.GetVectorTypeReference(Number<0>{})
+                                    .template AsType<AccDataType>()(t);
+                            c_acc_vec(t) = __builtin_elementwise_fma(
+                                c_partial_acc, c_scale_thread, c_acc_vec(t));
+                            anchor_accumulator_value(c_acc_vec(t));
                         });
                     });
 
@@ -606,14 +612,9 @@ struct BlockwiseGemmXdlops_pipeline_blockscale_bpreshuffle_v1<BlockGemmPipelineS
                     c_thread_buf_per_scale.GetVectorTypeReference(Number<0>{})
                         .template AsType<AccDataType>()(Number<t>{}) = 0;
                 });
-                vector_type<AccDataType, 2> c_scale_thread_vec;
                 constexpr index_t cscale_offset = CScaleThreadDesc{}.CalculateOffset(
                     make_tuple(kscale0, m0, n0 * num_scale_n_block / NRepeat));
-
-                c_scale_thread_vec.template AsType<AccDataType>()(Number<0>{}) =
-                    c_scale_thread_buf[Number<cscale_offset>{}];
-                c_scale_thread_vec.template AsType<AccDataType>()(Number<1>{}) =
-                    c_scale_thread_buf[Number<cscale_offset>{}];
+                const auto c_scale_thread = c_scale_thread_buf[Number<cscale_offset>{}];
 
                 static_for<0, KRepeat / num_scale_k_block, 1>{}([&](auto k0) {
                     vector_type<ComputeDataType, KPack> a_thread_vec;
@@ -653,16 +654,17 @@ struct BlockwiseGemmXdlops_pipeline_blockscale_bpreshuffle_v1<BlockGemmPipelineS
                 });
                 constexpr index_t c_offset = c_thread_desc_.CalculateOffset(make_tuple(m0, n0, 0));
 
-                static_for<0, xdlops_gemm.GetRegSizePerXdlops() / 2, 1>{}([&](auto t) {
-                    using pk_fma_type = typename vector_type<AccDataType, 2>::type;
-
-                    c_thread_buf.GetVectorTypeReference(Number<c_offset>{})
-                        .template AsType<pk_fma_type>()(t) = __builtin_elementwise_fma(
+                // Keep the post-scale FMA scalar; the old packed 2-lane update hits ROCm 7.2
+                // gfx950 illegal-type legalization/codegen issues.
+                static_for<0, xdlops_gemm.GetRegSizePerXdlops(), 1>{}([&](auto t) {
+                    auto& c_acc_vec = c_thread_buf.GetVectorTypeReference(Number<c_offset>{})
+                                          .template AsType<AccDataType>();
+                    const auto c_partial_acc =
                         c_thread_buf_per_scale.GetVectorTypeReference(Number<0>{})
-                            .template AsType<pk_fma_type>()[t],
-                        c_scale_thread_vec.template AsType<pk_fma_type>()[Number<0>{}],
-                        c_thread_buf.GetVectorTypeReference(Number<c_offset>{})
-                            .template AsType<pk_fma_type>()[t]);
+                            .template AsType<AccDataType>()(t);
+                    c_acc_vec(t) =
+                        __builtin_elementwise_fma(c_partial_acc, c_scale_thread, c_acc_vec(t));
+                    anchor_accumulator_value(c_acc_vec(t));
                 });
             });
 
@@ -701,14 +703,9 @@ struct BlockwiseGemmXdlops_pipeline_blockscale_bpreshuffle_v1<BlockGemmPipelineS
                     c_thread_buf_per_scale.GetVectorTypeReference(Number<0>{})
                         .template AsType<AccDataType>()(Number<t>{}) = 0;
                 });
-                vector_type<AccDataType, 2> c_scale_thread_vec;
                 constexpr index_t cscale_offset = CScaleThreadDesc{}.CalculateOffset(
                     make_tuple(kscale0, m0, n0 * num_scale_n_block / NRepeat));
-
-                c_scale_thread_vec.template AsType<AccDataType>()(Number<0>{}) =
-                    c_scale_thread_buf[Number<cscale_offset>{}];
-                c_scale_thread_vec.template AsType<AccDataType>()(Number<1>{}) =
-                    c_scale_thread_buf[Number<cscale_offset>{}];
+                const auto c_scale_thread = c_scale_thread_buf[Number<cscale_offset>{}];
 
                 static_for<0, KRepeat / num_scale_k_block, 1>{}([&](auto k0) {
                     vector_type<ComputeDataType, KPack> a_thread_vec;
@@ -748,16 +745,17 @@ struct BlockwiseGemmXdlops_pipeline_blockscale_bpreshuffle_v1<BlockGemmPipelineS
                 });
                 constexpr index_t c_offset = c_thread_desc_.CalculateOffset(make_tuple(m0, n0, 0));
 
-                static_for<0, xdlops_gemm.GetRegSizePerXdlops() / 2, 1>{}([&](auto t) {
-                    using pk_fma_type = typename vector_type<AccDataType, 2>::type;
-
-                    c_thread_buf.GetVectorTypeReference(Number<c_offset>{})
-                        .template AsType<pk_fma_type>()(t) = __builtin_elementwise_fma(
+                // Keep the post-scale FMA scalar; the old packed 2-lane update hits ROCm 7.2
+                // gfx950 illegal-type legalization/codegen issues.
+                static_for<0, xdlops_gemm.GetRegSizePerXdlops(), 1>{}([&](auto t) {
+                    auto& c_acc_vec = c_thread_buf.GetVectorTypeReference(Number<c_offset>{})
+                                          .template AsType<AccDataType>();
+                    const auto c_partial_acc =
                         c_thread_buf_per_scale.GetVectorTypeReference(Number<0>{})
-                            .template AsType<pk_fma_type>()[t],
-                        c_scale_thread_vec.template AsType<pk_fma_type>()[Number<0>{}],
-                        c_thread_buf.GetVectorTypeReference(Number<c_offset>{})
-                            .template AsType<pk_fma_type>()[t]);
+                            .template AsType<AccDataType>()(t);
+                    c_acc_vec(t) =
+                        __builtin_elementwise_fma(c_partial_acc, c_scale_thread, c_acc_vec(t));
+                    anchor_accumulator_value(c_acc_vec(t));
                 });
             });
         }
@@ -771,14 +769,9 @@ struct BlockwiseGemmXdlops_pipeline_blockscale_bpreshuffle_v1<BlockGemmPipelineS
                     c_thread_buf_per_scale.GetVectorTypeReference(Number<0>{})
                         .template AsType<AccDataType>()(Number<t>{}) = 0;
                 });
-                vector_type<AccDataType, 2> c_scale_thread_vec;
                 constexpr index_t cscale_offset = CScaleThreadDesc{}.CalculateOffset(
                     make_tuple(kscale0, m0, n0 * num_scale_n_block / NRepeat));
-
-                c_scale_thread_vec.template AsType<AccDataType>()(Number<0>{}) =
-                    c_scale_thread_buf[Number<cscale_offset>{}];
-                c_scale_thread_vec.template AsType<AccDataType>()(Number<1>{}) =
-                    c_scale_thread_buf[Number<cscale_offset>{}];
+                const auto c_scale_thread = c_scale_thread_buf[Number<cscale_offset>{}];
 
                 static_for<0, KRepeat / num_scale_k_block, 1>{}([&](auto k0) {
                     vector_type<ComputeDataType, KPack> a_thread_vec;
@@ -819,16 +812,17 @@ struct BlockwiseGemmXdlops_pipeline_blockscale_bpreshuffle_v1<BlockGemmPipelineS
 
                 constexpr index_t c_offset = c_thread_desc_.CalculateOffset(make_tuple(m0, n0, 0));
 
-                static_for<0, xdlops_gemm.GetRegSizePerXdlops() / 2, 1>{}([&](auto t) {
-                    using pk_fma_type = typename vector_type<AccDataType, 2>::type;
-
-                    c_thread_buf.GetVectorTypeReference(Number<c_offset>{})
-                        .template AsType<pk_fma_type>()(t) = __builtin_elementwise_fma(
+                // Keep the post-scale FMA scalar; the old packed 2-lane update hits ROCm 7.2
+                // gfx950 illegal-type legalization/codegen issues.
+                static_for<0, xdlops_gemm.GetRegSizePerXdlops(), 1>{}([&](auto t) {
+                    auto& c_acc_vec = c_thread_buf.GetVectorTypeReference(Number<c_offset>{})
+                                          .template AsType<AccDataType>();
+                    const auto c_partial_acc =
                         c_thread_buf_per_scale.GetVectorTypeReference(Number<0>{})
-                            .template AsType<pk_fma_type>()[t],
-                        c_scale_thread_vec.template AsType<pk_fma_type>()[Number<0>{}],
-                        c_thread_buf.GetVectorTypeReference(Number<c_offset>{})
-                            .template AsType<pk_fma_type>()[t]);
+                            .template AsType<AccDataType>()(t);
+                    c_acc_vec(t) =
+                        __builtin_elementwise_fma(c_partial_acc, c_scale_thread, c_acc_vec(t));
+                    anchor_accumulator_value(c_acc_vec(t));
                 });
             });
         }
diff --git a/projects/composablekernel/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_blockscale_b_preshuffle_v3.hpp b/projects/composablekernel/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_blockscale_b_preshuffle_v3.hpp
index da260fb5a10c..739b3152a322 100644
--- a/projects/composablekernel/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_blockscale_b_preshuffle_v3.hpp
+++ b/projects/composablekernel/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_blockscale_b_preshuffle_v3.hpp
@@ -416,6 +416,15 @@ struct BlockwiseGemmXdlops_pipeline_blockscale_bpreshuffle_v3<BlockGemmPipelineS
     {
         ignore = b_block_desc;
         ignore = b_block_buf;
+
+        static_assert(sizeof(AccDataType) == 4 && is_floating_point<AccDataType>::value,
+                      "bpreshuffle accumulator anchor requires 32-bit float AccDataType");
+
+        // The empty asm is a read/write VGPR use of the updated accumulator.
+        // It keeps each post-scale accumulator definition visible to the
+        // optimizer, which enables determinism for bitwise-stable repeated launches.
+        auto anchor_accumulator_value = [&](auto& value) { asm volatile("" : "+v"(value)); };
+
         __builtin_amdgcn_sched_barrier(0);
         static_assert(CScaleThreadDesc{}.GetLength(Number<0>{}) == 1,
                       "Pipeline v3 only support scaleblocksliceK=1");
@@ -642,11 +651,7 @@ struct BlockwiseGemmXdlops_pipeline_blockscale_bpreshuffle_v3<BlockGemmPipelineS
                     }
 
                     static_for<0, MRepeat, 1>{}([&](auto m0) {
-                        vector_type<AccDataType, 2> c_scale_thread_vec;
-                        c_scale_thread_vec.template AsType<AccDataType>()(Number<0>{}) =
-                            c_scale_thread_buf[m0];
-                        c_scale_thread_vec.template AsType<AccDataType>()(Number<1>{}) =
-                            c_scale_thread_buf[m0];
+                        const auto c_scale_thread = c_scale_thread_buf[m0];
 
                         static_for<0, NRepeat, 1>{}([&](auto n0) {
                             constexpr auto mfma_buf_offset =
@@ -712,17 +717,19 @@ struct BlockwiseGemmXdlops_pipeline_blockscale_bpreshuffle_v3<BlockGemmPipelineS
                             constexpr index_t c_offset =
                                 c_thread_desc_.CalculateOffset(make_tuple(m0, n0, 0));
 
-                            static_for<0, xdlops_gemm.GetRegSizePerXdlops() / 2, 1>{}([&](auto t) {
-                                using pk_fma_type = typename vector_type<AccDataType, 2>::type;
-
-                                c_thread_buf.GetVectorTypeReference(Number<c_offset>{})
-                                    .template AsType<pk_fma_type>()(t) = __builtin_elementwise_fma(
+                            // Keep the post-scale FMA scalar; the old packed 2-lane update hits
+                            // ROCm 7.2 gfx950 illegal-type legalization/codegen issues.
+                            static_for<0, xdlops_gemm.GetRegSizePerXdlops(), 1>{}([&](auto t) {
+                                auto& c_acc_vec =
+                                    c_thread_buf.GetVectorTypeReference(Number<c_offset>{})
+                                        .template AsType<AccDataType>();
+                                const auto c_partial_acc =
                                     c_thread_buf_per_scale
                                         .GetVectorTypeReference(Number<scale_buf_offset>{})
-                                        .template AsType<pk_fma_type>()[t],
-                                    c_scale_thread_vec.template AsType<pk_fma_type>()[Number<0>{}],
-                                    c_thread_buf.GetVectorTypeReference(Number<c_offset>{})
-                                        .template AsType<pk_fma_type>()[t]);
+                                        .template AsType<AccDataType>()(t);
+                                c_acc_vec(t) = __builtin_elementwise_fma(
+                                    c_partial_acc, c_scale_thread, c_acc_vec(t));
+                                anchor_accumulator_value(c_acc_vec(t));
                             });
                         });
 
@@ -793,11 +800,7 @@ struct BlockwiseGemmXdlops_pipeline_blockscale_bpreshuffle_v3<BlockGemmPipelineS
             a_blockwise_copy.RunWrite(a_block_desc, a_block_buf.At(I1));
 
             static_for<0, MRepeat, 1>{}([&](auto m0) {
-                vector_type<AccDataType, 2> c_scale_thread_vec;
-                c_scale_thread_vec.template AsType<AccDataType>()(Number<0>{}) =
-                    c_scale_thread_buf[m0];
-                c_scale_thread_vec.template AsType<AccDataType>()(Number<1>{}) =
-                    c_scale_thread_buf[m0];
+                const auto c_scale_thread = c_scale_thread_buf[m0];
 
                 static_for<0, NRepeat, 1>{}([&](auto n0) {
                     constexpr auto mfma_buf_offset =
@@ -858,17 +861,18 @@ struct BlockwiseGemmXdlops_pipeline_blockscale_bpreshuffle_v3<BlockGemmPipelineS
                     constexpr index_t c_offset =
                         c_thread_desc_.CalculateOffset(make_tuple(m0, n0, 0));
 
-                    static_for<0, xdlops_gemm.GetRegSizePerXdlops() / 2, 1>{}([&](auto t) {
-                        using pk_fma_type = typename vector_type<AccDataType, 2>::type;
-
-                        c_thread_buf.GetVectorTypeReference(Number<c_offset>{})
-                            .template AsType<pk_fma_type>()(t) = __builtin_elementwise_fma(
+                    // Keep the post-scale FMA scalar; the old packed 2-lane update hits ROCm 7.2
+                    // gfx950 illegal-type legalization/codegen issues.
+                    static_for<0, xdlops_gemm.GetRegSizePerXdlops(), 1>{}([&](auto t) {
+                        auto& c_acc_vec = c_thread_buf.GetVectorTypeReference(Number<c_offset>{})
+                                              .template AsType<AccDataType>();
+                        const auto c_partial_acc =
                             c_thread_buf_per_scale
                                 .GetVectorTypeReference(Number<scale_buf_offset>{})
-                                .template AsType<pk_fma_type>()[t],
-                            c_scale_thread_vec.template AsType<pk_fma_type>()[Number<0>{}],
-                            c_thread_buf.GetVectorTypeReference(Number<c_offset>{})
-                                .template AsType<pk_fma_type>()[t]);
+                                .template AsType<AccDataType>()(t);
+                        c_acc_vec(t) =
+                            __builtin_elementwise_fma(c_partial_acc, c_scale_thread, c_acc_vec(t));
+                        anchor_accumulator_value(c_acc_vec(t));
                     });
                 });
 
@@ -909,11 +913,7 @@ struct BlockwiseGemmXdlops_pipeline_blockscale_bpreshuffle_v3<BlockGemmPipelineS
             });
 
             static_for<0, MRepeat, 1>{}([&](auto m0) {
-                vector_type<AccDataType, 2> c_scale_thread_vec;
-                c_scale_thread_vec.template AsType<AccDataType>()(Number<0>{}) =
-                    c_scale_thread_buf[m0];
-                c_scale_thread_vec.template AsType<AccDataType>()(Number<1>{}) =
-                    c_scale_thread_buf[m0];
+                const auto c_scale_thread = c_scale_thread_buf[m0];
 
                 static_for<0, NRepeat, 1>{}([&](auto n0) {
                     constexpr auto mfma_buf_offset =
@@ -974,17 +974,18 @@ struct BlockwiseGemmXdlops_pipeline_blockscale_bpreshuffle_v3<BlockGemmPipelineS
                     constexpr index_t c_offset =
                         c_thread_desc_.CalculateOffset(make_tuple(m0, n0, 0));
 
-                    static_for<0, xdlops_gemm.GetRegSizePerXdlops() / 2, 1>{}([&](auto t) {
-                        using pk_fma_type = typename vector_type<AccDataType, 2>::type;
-
-                        c_thread_buf.GetVectorTypeReference(Number<c_offset>{})
-                            .template AsType<pk_fma_type>()(t) = __builtin_elementwise_fma(
+                    // Keep the post-scale FMA scalar; the old packed 2-lane update hits ROCm 7.2
+                    // gfx950 illegal-type legalization/codegen issues.
+                    static_for<0, xdlops_gemm.GetRegSizePerXdlops(), 1>{}([&](auto t) {
+                        auto& c_acc_vec = c_thread_buf.GetVectorTypeReference(Number<c_offset>{})
+                                              .template AsType<AccDataType>();
+                        const auto c_partial_acc =
                             c_thread_buf_per_scale
                                 .GetVectorTypeReference(Number<scale_buf_offset>{})
-                                .template AsType<pk_fma_type>()[t],
-                            c_scale_thread_vec.template AsType<pk_fma_type>()[Number<0>{}],
-                            c_thread_buf.GetVectorTypeReference(Number<c_offset>{})
-                                .template AsType<pk_fma_type>()[t]);
+                                .template AsType<AccDataType>()(t);
+                        c_acc_vec(t) =
+                            __builtin_elementwise_fma(c_partial_acc, c_scale_thread, c_acc_vec(t));
+                        anchor_accumulator_value(c_acc_vec(t));
                     });
                 });
 
@@ -1021,11 +1022,7 @@ struct BlockwiseGemmXdlops_pipeline_blockscale_bpreshuffle_v3<BlockGemmPipelineS
         else
         {
             static_for<0, MRepeat, 1>{}([&](auto m0) {
-                vector_type<AccDataType, 2> c_scale_thread_vec;
-                c_scale_thread_vec.template AsType<AccDataType>()(Number<0>{}) =
-                    c_scale_thread_buf[m0];
-                c_scale_thread_vec.template AsType<AccDataType>()(Number<1>{}) =
-                    c_scale_thread_buf[m0];
+                const auto c_scale_thread = c_scale_thread_buf[m0];
 
                 static_for<0, NRepeat, 1>{}([&](auto n0) {
                     constexpr auto mfma_buf_offset =
@@ -1087,17 +1084,18 @@ struct BlockwiseGemmXdlops_pipeline_blockscale_bpreshuffle_v3<BlockGemmPipelineS
                     constexpr index_t c_offset =
                         c_thread_desc_.CalculateOffset(make_tuple(m0, n0, 0));
 
-                    static_for<0, xdlops_gemm.GetRegSizePerXdlops() / 2, 1>{}([&](auto t) {
-                        using pk_fma_type = typename vector_type<AccDataType, 2>::type;
-
-                        c_thread_buf.GetVectorTypeReference(Number<c_offset>{})
-                            .template AsType<pk_fma_type>()(t) = __builtin_elementwise_fma(
+                    // Keep the post-scale FMA scalar; the old packed 2-lane update hits ROCm 7.2
+                    // gfx950 illegal-type legalization/codegen issues.
+                    static_for<0, xdlops_gemm.GetRegSizePerXdlops(), 1>{}([&](auto t) {
+                        auto& c_acc_vec = c_thread_buf.GetVectorTypeReference(Number<c_offset>{})
+                                              .template AsType<AccDataType>();
+                        const auto c_partial_acc =
                             c_thread_buf_per_scale
                                 .GetVectorTypeReference(Number<scale_buf_offset>{})
-                                .template AsType<pk_fma_type>()[t],
-                            c_scale_thread_vec.template AsType<pk_fma_type>()[Number<0>{}],
-                            c_thread_buf.GetVectorTypeReference(Number<c_offset>{})
-                                .template AsType<pk_fma_type>()[t]);
+                                .template AsType<AccDataType>()(t);
+                        c_acc_vec(t) =
+                            __builtin_elementwise_fma(c_partial_acc, c_scale_thread, c_acc_vec(t));
+                        anchor_accumulator_value(c_acc_vec(t));
                     });
                 });
 
diff --git a/projects/composablekernel/profiler/include/profiler/profile_gemm_blockscale_wp_impl.hpp b/projects/composablekernel/profiler/include/profiler/profile_gemm_blockscale_wp_impl.hpp
index 9e5694b71146..62daa69309a7 100644
--- a/projects/composablekernel/profiler/include/profiler/profile_gemm_blockscale_wp_impl.hpp
+++ b/projects/composablekernel/profiler/include/profiler/profile_gemm_blockscale_wp_impl.hpp
@@ -3,8 +3,13 @@
 
 #pragma once
 
+#include <algorithm>
+#include <cmath>
+#include <cstring>
 #include <iomanip>
 #include <iostream>
+#include <memory>
+#include <string>
 #include <typeinfo>
 
 #include "ck/ck.hpp"
@@ -69,6 +74,7 @@ template <typename A0DataType,
           typename ALayout,
           typename BLayout,
           typename ELayout>
+// clang-format off
 bool profile_gemm_blockscale_weightpreshuffle_impl(int do_verification,
                                                    int init_method,
                                                    bool do_log,
@@ -81,11 +87,15 @@ bool profile_gemm_blockscale_weightpreshuffle_impl(int do_verification,
                                                    int StrideE,
                                                    int n_warmup,
                                                    int n_iter,
-                                                   uint64_t rotating  = 0,
+                                                   uint64_t rotating = 0,
+                                                   int determinism_check = 1,
                                                    int instance_index = -1)
+// clang-format on
 {
     bool pass = true;
 
+    determinism_check = std::max(1, determinism_check);
+
     auto f_host_tensor_descriptor = [](std::size_t row, std::size_t col, int& stride, auto layout) {
         using namespace ck::literals;
 
@@ -228,6 +238,71 @@ bool profile_gemm_blockscale_weightpreshuffle_impl(int do_verification,
 
     std::cout << "found " << op_ptrs.size() << " instances" << std::endl;
 
+    auto check_row_relative_error =
+        [&](const auto& actual, const auto& expected, const std::string& label) {
+            // sglang#28685 showed sparse per-row spikes near tile boundaries.
+            constexpr double kReferenceDenominatorFloor = 1.0;
+            // Loose enough for expected FP8 noise, tight enough to catch the
+            // >10% per-row spikes reported for the miscompiled gfx950 path.
+            constexpr double kBadRowRelativeThreshold = 0.1;
+            constexpr int kBadRowSampleLimit          = 16;
+
+            double worst_row_rel_max_error = -1;
+            int worst_row_rel              = 0;
+            int bad_row_count              = 0;
+            int sampled_bad_row_count      = 0;
+            std::string bad_row_sample;
+
+            for(int m = 0; m < M; ++m)
+            {
+                double row_max_abs_error = 0;
+                double row_ref_max_abs   = 0;
+
+                for(int n = 0; n < N; ++n)
+                {
+                    const float actual_value   = ck::type_convert<float>(actual(m, n));
+                    const float expected_value = ck::type_convert<float>(expected(m, n));
+                    row_max_abs_error          = std::max(row_max_abs_error,
+                                                 std::abs(static_cast<double>(actual_value) -
+                                                          static_cast<double>(expected_value)));
+                    row_ref_max_abs =
+                        std::max(row_ref_max_abs, std::abs(static_cast<double>(expected_value)));
+                }
+
+                const double row_rel_max_error =
+                    row_max_abs_error / std::max(row_ref_max_abs, kReferenceDenominatorFloor);
+                if(row_rel_max_error > worst_row_rel_max_error)
+                {
+                    worst_row_rel_max_error = row_rel_max_error;
+                    worst_row_rel           = m;
+                }
+                if(row_rel_max_error > kBadRowRelativeThreshold)
+                {
+                    ++bad_row_count;
+                    if(sampled_bad_row_count < kBadRowSampleLimit)
+                    {
+                        if(!bad_row_sample.empty())
+                        {
+                            bad_row_sample += ",";
+                        }
+                        bad_row_sample += std::to_string(m);
+                        ++sampled_bad_row_count;
+                    }
+                }
+            }
+
+            const bool pass_row_relative_check = bad_row_count == 0;
+            if(!pass_row_relative_check)
+            {
+                std::cout << label << " rowrel_max=" << worst_row_rel_max_error
+                          << " rowrel_worst_row=" << worst_row_rel
+                          << " rowrel_bad_rows_gt_0p1=" << bad_row_count << "/" << M
+                          << " rowrel_bad_row_sample=" << bad_row_sample << std::endl;
+            }
+
+            return pass_row_relative_check;
+        };
+
     // Run reference GEMM
     if(do_verification)
     {
@@ -318,43 +393,58 @@ bool profile_gemm_blockscale_weightpreshuffle_impl(int do_verification,
 
         if(op_ptr->IsSupportedArgument(argument_ptr.get()))
         {
-
             // re-init C to zero before profiling next kernel
             c_device_buf.SetZero();
 
             invoker_ptr->Run(argument_ptr.get(), StreamConfig{nullptr, false, 0, n_warmup, n_iter});
 
-            if(do_verification)
+            std::unique_ptr<Tensor<EDataType>> first_device_result;
+            if(determinism_check > 1)
+            {
+                first_device_result =
+                    std::make_unique<Tensor<EDataType>>(e_m_n_device_result.mDesc);
+            }
+
+            if(do_verification || determinism_check > 1)
             {
                 c_device_buf.FromDevice(e_m_n_device_result.mData.data());
+                if(first_device_result)
+                {
+                    first_device_result->mData = e_m_n_device_result.mData;
+                }
+            }
 
+            if(do_verification)
+            {
+                bool current_pass = true;
 #if defined CK_ENABLE_FP8
                 // set softer tolerances for fp8
                 if constexpr(is_same_v<A0DataType, f8_t> || is_same_v<B0DataType, f8_t> ||
                              is_same_v<EDataType, f8_t>)
                 {
-                    std::string msg   = "Error: Incorrect results!";
-                    double rtol       = 5e-2;
-                    double atol       = 5e-2;
-                    bool current_pass = ck::utils::check_err(
+                    std::string msg = "Error: Incorrect results!";
+                    double rtol     = 5e-2;
+                    double atol     = 5e-2;
+                    current_pass    = ck::utils::check_err(
                         e_m_n_device_result, e_m_n_host_result, msg, rtol, atol);
-                    pass = pass & current_pass;
-                    if(!current_pass)
-                    {
-                        std::cout << op_ptr->GetTypeString() << " failed" << std::endl;
-                    }
                 }
                 else
                 {
 #endif
-                    pass = pass & ck::utils::check_err(e_m_n_device_result, e_m_n_host_result);
-                    if(!pass)
-                    {
-                        std::cout << op_ptr->GetTypeString() << " failed" << std::endl;
-                    }
+                    current_pass = ck::utils::check_err(e_m_n_device_result, e_m_n_host_result);
 #if defined CK_ENABLE_FP8
                 }
 #endif
+                // Keep sparse per-row coverage independent of the dtype tolerance branch.
+                current_pass = check_row_relative_error(e_m_n_device_result,
+                                                        e_m_n_host_result,
+                                                        "Reference row-relative check") &&
+                               current_pass;
+                pass = pass & current_pass;
+                if(!current_pass)
+                {
+                    std::cout << op_ptr->GetTypeString() << " failed" << std::endl;
+                }
 
                 if(do_log)
                 {
@@ -367,6 +457,31 @@ bool profile_gemm_blockscale_weightpreshuffle_impl(int do_verification,
                 }
             }
 
+            for(int repeat = 1; repeat < determinism_check; ++repeat)
+            {
+                // These extra single-launch repeats sample the fixed-shape instability.
+                // Use a bit-exact memcmp; a tolerance compare would hide hash drift.
+                c_device_buf.SetZero();
+                invoker_ptr->Run(argument_ptr.get(), StreamConfig{nullptr, false, 0, 0, 1});
+                c_device_buf.FromDevice(e_m_n_device_result.mData.data());
+
+                const auto byte_count    = first_device_result->mData.size() * sizeof(EDataType);
+                const bool deterministic = std::memcmp(e_m_n_device_result.mData.data(),
+                                                       first_device_result->mData.data(),
+                                                       byte_count) == 0;
+
+                if(!deterministic)
+                {
+                    pass = false;
+                    std::cout << op_ptr->GetTypeString()
+                              << " produced nondeterministic output on repeat " << repeat << " of "
+                              << determinism_check << std::endl;
+                    check_row_relative_error(
+                        e_m_n_device_result, *first_device_result, "Determinism mismatch");
+                    break;
+                }
+            }
+
             std::string op_name = op_ptr->GetTypeString();
 
             float ave_time = invoker_ptr->Run(
diff --git a/projects/composablekernel/test/gemm_blockscale_wp/test_gemm_blockscale_wp_fp8.cpp b/projects/composablekernel/test/gemm_blockscale_wp/test_gemm_blockscale_wp_fp8.cpp
index 71ab22b77db7..0533fd6235df 100644
--- a/projects/composablekernel/test/gemm_blockscale_wp/test_gemm_blockscale_wp_fp8.cpp
+++ b/projects/composablekernel/test/gemm_blockscale_wp/test_gemm_blockscale_wp_fp8.cpp
@@ -1,6 +1,8 @@
 // Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
 
+#include <cstdlib>
+#include <cstring>
 #include <tuple>
 
 #include "gtest/gtest.h"
@@ -25,6 +27,12 @@ struct tuple_concat<std::tuple<Xs...>, std::tuple<Ys...>>
     using type = std::tuple<Xs..., Ys...>;
 };
 
+bool run_long_gemm_regression_tests()
+{
+    const char* value = std::getenv("CK_RUN_LONG_GEMM_REGRESSION_TESTS");
+    return value != nullptr && std::strcmp(value, "0") != 0;
+}
+
 } // namespace
 
 template <typename Tuple>
@@ -50,5 +58,68 @@ TYPED_TEST(TestGemmBlockScaleWP_FP8_MK_NK, Regular0)
     constexpr int K = 512;
 
     for(int M : Ms)
+    {
+        SCOPED_TRACE(::testing::Message() << "M=" << M << " N=" << N << " K=" << K);
         this->Run(M, N, K);
+    }
+}
+
+TYPED_TEST(TestGemmBlockScaleWP_FP8_MK_NK, ReportedWkvDeterminism)
+{
+    if(!run_long_gemm_regression_tests())
+    {
+        GTEST_SKIP() << "Set CK_RUN_LONG_GEMM_REGRESSION_TESTS=1 to run the ROCm/aiter#3261 "
+                        "8192x512xK determinism repro shapes.";
+    }
+
+    // Large WKV shapes from ROCm/aiter#3261. Check determinism only.
+    constexpr int M    = 8192;
+    constexpr int N    = 512;
+    constexpr int Ks[] = {384, 640, 3968, 4096, 4224};
+
+    for(const int K : Ks)
+    {
+        SCOPED_TRACE(::testing::Message() << "M=" << M << " N=" << N << " K=" << K);
+        // Determinism-only large shapes; the next test covers host-reference accuracy.
+        this->Run(M, N, K, 0, 1, 8, false);
+    }
+}
+
+TYPED_TEST(TestGemmBlockScaleWP_FP8_MK_NK, Glm52OutOfAllowlistAccuracyAndDeterminism)
+{
+    if(!run_long_gemm_regression_tests())
+    {
+        GTEST_SKIP() << "Set CK_RUN_LONG_GEMM_REGRESSION_TESTS=1 to run the GLM-5.2-FP8 "
+                        "out-of-allowlist bpreshuffle accuracy and determinism regression shapes.";
+    }
+
+    constexpr std::tuple<int, int, int> shapes[] = {
+        // GLM-5.2-FP8 projection shapes from sgl-project/sglang#28685.
+        // The kv_b_proj cases cover the reported M-row boundary failures.
+        {8, 2048, 6144},   // q_a_proj
+        {8, 16384, 2048},  // q_b_proj
+        {64, 16384, 2048}, // q_b_proj
+        {12, 6144, 12288}, // mlp.down_proj
+        {64, 6144, 12288}, // mlp.down_proj
+        {8, 28672, 512},   // kv_b_proj
+        {12, 28672, 512},  // kv_b_proj
+        {32, 28672, 512},  // kv_b_proj
+        {48, 28672, 512},  // kv_b_proj
+        {56, 28672, 512},  // kv_b_proj
+        {64, 28672, 512},  // kv_b_proj
+        {72, 28672, 512},  // kv_b_proj
+        {96, 28672, 512},  // kv_b_proj
+        {128, 28672, 512}, // kv_b_proj
+    };
+
+    for(const auto& shape : shapes)
+    {
+        const int M = std::get<0>(shape);
+        const int N = std::get<1>(shape);
+        const int K = std::get<2>(shape);
+        SCOPED_TRACE(::testing::Message() << "M=" << M << " N=" << N << " K=" << K);
+
+        // Compare to host reference and repeat the device run.
+        this->Run(M, N, K, 0, 1, 4, true);
+    }
 }
diff --git a/projects/composablekernel/test/gemm_blockscale_wp/test_gemm_common.hpp b/projects/composablekernel/test/gemm_blockscale_wp/test_gemm_common.hpp
index 34cf330519e4..f88f06a530d8 100644
--- a/projects/composablekernel/test/gemm_blockscale_wp/test_gemm_common.hpp
+++ b/projects/composablekernel/test/gemm_blockscale_wp/test_gemm_common.hpp
@@ -37,7 +37,13 @@ class TestGemmBlockscaleWPCommon : public ::testing::Test
     static constexpr index_t ScaleBlockN = 128;
     static constexpr index_t ScaleBlockK = 128;
 
-    void Run(const int M, const int N, const int K, int n_warmup = 1, int n_iter = 10)
+    void Run(const int M,
+             const int N,
+             const int K,
+             int n_warmup          = 1,
+             int n_iter            = 10,
+             int determinism_check = 1,
+             bool do_verification  = verify_)
     {
         bool all_success = true;
 
@@ -59,7 +65,7 @@ class TestGemmBlockscaleWPCommon : public ::testing::Test
                                                                         ScaleBlockK,
                                                                         ALayout,
                                                                         BLayout,
-                                                                        CLayout>(verify_,
+                                                                        CLayout>(do_verification,
                                                                                  init_method_,
                                                                                  log_,
                                                                                  bench_,
@@ -72,6 +78,7 @@ class TestGemmBlockscaleWPCommon : public ::testing::Test
                                                                                  n_warmup,
                                                                                  n_iter,
                                                                                  0,
+                                                                                 determinism_check,
                                                                                  instance_index);
 
         EXPECT_TRUE(all_success);