compile time fix for gfx1201

Author: amd-khushbu

include/ck_tile/ops/gemm_quant/pipeline/gemm_quant_pipeline_problem.hpp

@@ -79,10 +79,8 @@ struct GemmQuantPipelineProblemBase : public GemmPipelineProblemBase<ADataType_,
     static constexpr auto TailNum    = TailNum_;
 
     static_assert(BlockGemmShape::kM % AQuantGroupSize::kM == 0);
-    static_assert(BlockGemmShape::kN % AQuantGroupSize::kN == 0);
     static_assert(BlockGemmShape::kK % AQuantGroupSize::kK == 0);
     static_assert(BlockGemmShape::kM % BQuantGroupSize::kM == 0);
-    static_assert(BlockGemmShape::kN % BQuantGroupSize::kN == 0);
     static_assert(BlockGemmShape::kK % BQuantGroupSize::kK == 0);
 
     [[nodiscard]] CK_TILE_HOST static const std::string GetName()

include/ck_tile/ops/gemm_quant/pipeline/gemm_wp_bquant_pipeline_ag_bg_cr_v2.hpp

@@ -144,23 +144,32 @@ struct WPQuantBPipelineAgBgCrV2 : public WeightPreshufflePipelineAGmemBGmemCRegV
                 // Insert LDS read/write groups periodically based on ds_rep.
                 // The % pattern staggers READ and WRITE so they don't collapse
                 // into the same cycle in the model.
-                if constexpr(ds_rep > 0 && i_inst % ds_rep == 0)
+                if constexpr(ds_rep > 0)
                 {
-                    __builtin_amdgcn_sched_group_barrier(
-                        LLVMSchedGroupMask::DS_READ, 1, 0); // DS read
+                    if( i_inst % ds_rep == 0)
+                    {
+                        __builtin_amdgcn_sched_group_barrier(
+                            LLVMSchedGroupMask::DS_READ, 1, 0); // DS read
+                    }
                 }
-                if constexpr(ds_rep > 0 && i_inst % ds_rep == 1)
+                if constexpr(ds_rep > 0)
                 {
-                    __builtin_amdgcn_sched_group_barrier(
-                        LLVMSchedGroupMask::DS_WRITE, 1, 0); // DS write
+                    if(i_inst % ds_rep == 1)
+                    {
+                        __builtin_amdgcn_sched_group_barrier(
+                            LLVMSchedGroupMask::DS_WRITE, 1, 0); // DS write
+                    }
                 }
 
-                if constexpr(buffer_load_rep > 0 && i_inst % buffer_load_rep == 0)
+                if constexpr(buffer_load_rep > 0)
                 {
-                    if constexpr(ds_write_inst > 0)
+                    if(i_inst % buffer_load_rep == 0)
                     {
-                        __builtin_amdgcn_sched_group_barrier(
-                            LLVMSchedGroupMask::VMEM_READ, 1, 0); // VMEM read
+                        if constexpr(ds_write_inst > 0)
+                        {
+                            __builtin_amdgcn_sched_group_barrier(
+                                LLVMSchedGroupMask::VMEM_READ, 1, 0); // VMEM read
+                        }
                     }
                 }
                 // Always mark some VALU work in the loop to reflect auxiliary scalar