src/amd/compiler/aco_assembler.cpp

@@ -1537,7 +1537,8 @@ emit_long_jump(asm_context& ctx, SALU_instruction* branch, bool backwards,
       case aco_opcode::s_cbranch_execnz: inv = aco_opcode::s_cbranch_execz; break;
       default: unreachable("Unhandled long jump.");
       }
-      instr.reset(bld.sopp(inv, 6));
+      unsigned size = ctx.gfx_level >= GFX12 ? 7 : 6;
+      instr.reset(bld.sopp(inv, size));
       emit_sopp_instruction(ctx, out, instr.get(), true);
    }
 
@@ -1545,6 +1546,11 @@ emit_long_jump(asm_context& ctx, SALU_instruction* branch, bool backwards,
    instr.reset(bld.sop1(aco_opcode::s_getpc_b64, def).instr);
    emit_instruction(ctx, out, instr.get());
 
+   if (ctx.gfx_level >= GFX12) {
+      instr.reset(bld.sop1(aco_opcode::s_sext_i32_i16, def_tmp_hi, op_tmp_hi).instr);
+      emit_instruction(ctx, out, instr.get());
+   }
+
    instr.reset(
       bld.sop2(aco_opcode::s_addc_u32, def_tmp_lo, op_tmp_lo, Operand::literal32(0)).instr);
    emit_instruction(ctx, out, instr.get());

src/amd/compiler/aco_insert_NOPs.cpp

@@ -253,6 +253,8 @@ struct NOP_ctx_gfx11 {
 
    /* LdsDirectVMEMHazard */
    std::bitset<256> vgpr_used_by_vmem_load;
+   std::bitset<256> vgpr_used_by_vmem_sample;
+   std::bitset<256> vgpr_used_by_vmem_bvh;
    std::bitset<256> vgpr_used_by_vmem_store;
    std::bitset<256> vgpr_used_by_ds;
 
@@ -268,6 +270,8 @@ struct NOP_ctx_gfx11 {
    {
       has_Vcmpx |= other.has_Vcmpx;
       vgpr_used_by_vmem_load |= other.vgpr_used_by_vmem_load;
+      vgpr_used_by_vmem_sample |= other.vgpr_used_by_vmem_sample;
+      vgpr_used_by_vmem_bvh |= other.vgpr_used_by_vmem_bvh;
       vgpr_used_by_vmem_store |= other.vgpr_used_by_vmem_store;
       vgpr_used_by_ds |= other.vgpr_used_by_ds;
       valu_since_wr_by_trans.join_min(other.valu_since_wr_by_trans);
@@ -281,6 +285,8 @@ struct NOP_ctx_gfx11 {
    {
       return has_Vcmpx == other.has_Vcmpx &&
              vgpr_used_by_vmem_load == other.vgpr_used_by_vmem_load &&
+             vgpr_used_by_vmem_sample == other.vgpr_used_by_vmem_sample &&
+             vgpr_used_by_vmem_bvh == other.vgpr_used_by_vmem_bvh &&
              vgpr_used_by_vmem_store == other.vgpr_used_by_vmem_store &&
              vgpr_used_by_ds == other.vgpr_used_by_ds &&
              valu_since_wr_by_trans == other.valu_since_wr_by_trans &&
@@ -1373,7 +1379,9 @@ handle_instruction_gfx11(State& state, NOP_ctx_gfx11& ctx, aco_ptr<Instruction>&
       ctx.has_Vcmpx = true;
    } else if (ctx.has_Vcmpx && (instr->opcode == aco_opcode::v_permlane16_b32 ||
                                 instr->opcode == aco_opcode::v_permlanex16_b32 ||
-                                instr->opcode == aco_opcode::v_permlane64_b32)) {
+                                instr->opcode == aco_opcode::v_permlane64_b32 ||
+                                instr->opcode == aco_opcode::v_permlane16_var_b32 ||
+                                instr->opcode == aco_opcode::v_permlanex16_var_b32)) {
       ctx.has_Vcmpx = false;
 
       /* Unlike on GFX10, v_nop should resolve the hazard on GFX11. */
@@ -1395,6 +1403,8 @@ handle_instruction_gfx11(State& state, NOP_ctx_gfx11& ctx, aco_ptr<Instruction>&
       /* va_vdst already obtained through parse_vdst_wait(). */
       vm_vsrc = (instr->salu().imm >> 2) & 0x7;
       sa_sdst = instr->salu().imm & 0x1;
+   } else if (instr->isLDSDIR() && state.program->gfx_level >= GFX12) {
+      vm_vsrc = instr->ldsdir().wait_vsrc ? 7 : 0;
    }
 
    if (instr->isLDSDIR()) {
@@ -1410,7 +1420,7 @@ handle_instruction_gfx11(State& state, NOP_ctx_gfx11& ctx, aco_ptr<Instruction>&
     * VALU reads VGPR written by transcendental instruction without 6+ VALU or 2+ transcendental
     * in-between.
     */
-   if (va_vdst > 0 && instr->isVALU()) {
+   if (state.program->gfx_level < GFX11_5 && va_vdst > 0 && instr->isVALU()) {
       uint8_t num_valu = 15;
       uint8_t num_trans = 15;
       for (Operand& op : instr->operands) {
@@ -1427,65 +1437,68 @@ handle_instruction_gfx11(State& state, NOP_ctx_gfx11& ctx, aco_ptr<Instruction>&
       }
    }
 
-   if (va_vdst > 0 && handle_valu_partial_forwarding_hazard(state, instr)) {
+   if (va_vdst > 0 && state.program->gfx_level < GFX12 &&
+       handle_valu_partial_forwarding_hazard(state, instr)) {
       bld.sopp(aco_opcode::s_waitcnt_depctr, 0x0fff);
       va_vdst = 0;
    }
 
-   /* VALUMaskWriteHazard
-    * VALU reads SGPR as a lane mask and later written by SALU cannot safely be read by SALU.
-    */
-   if (state.program->wave_size == 64 && instr->isSALU() &&
-       check_written_regs(instr, ctx.sgpr_read_by_valu_as_lanemask)) {
-      ctx.sgpr_read_by_valu_as_lanemask_then_wr_by_salu = ctx.sgpr_read_by_valu_as_lanemask;
-      ctx.sgpr_read_by_valu_as_lanemask.reset();
-   } else if (state.program->wave_size == 64 && instr->isSALU() &&
-              check_read_regs(instr, ctx.sgpr_read_by_valu_as_lanemask_then_wr_by_salu)) {
-      bld.sopp(aco_opcode::s_waitcnt_depctr, 0xfffe);
-      sa_sdst = 0;
-   }
+   if (state.program->gfx_level < GFX12) {
+      /* VALUMaskWriteHazard
+       * VALU reads SGPR as a lane mask and later written by SALU cannot safely be read by SALU.
+       */
+      if (state.program->wave_size == 64 && instr->isSALU() &&
+          check_written_regs(instr, ctx.sgpr_read_by_valu_as_lanemask)) {
+         ctx.sgpr_read_by_valu_as_lanemask_then_wr_by_salu = ctx.sgpr_read_by_valu_as_lanemask;
+         ctx.sgpr_read_by_valu_as_lanemask.reset();
+      } else if (state.program->wave_size == 64 && instr->isSALU() &&
+                 check_read_regs(instr, ctx.sgpr_read_by_valu_as_lanemask_then_wr_by_salu)) {
+         bld.sopp(aco_opcode::s_waitcnt_depctr, 0xfffe);
+         sa_sdst = 0;
+      }
 
-   if (va_vdst == 0) {
-      ctx.valu_since_wr_by_trans.reset();
-      ctx.trans_since_wr_by_trans.reset();
-   }
+      if (va_vdst == 0) {
+         ctx.valu_since_wr_by_trans.reset();
+         ctx.trans_since_wr_by_trans.reset();
+      }
 
-   if (sa_sdst == 0)
-      ctx.sgpr_read_by_valu_as_lanemask_then_wr_by_salu.reset();
+      if (sa_sdst == 0)
+         ctx.sgpr_read_by_valu_as_lanemask_then_wr_by_salu.reset();
 
-   if (instr->isVALU()) {
-      bool is_trans = instr->isTrans();
+      if (instr->isVALU()) {
+         bool is_trans = instr->isTrans();
 
-      ctx.valu_since_wr_by_trans.inc();
-      if (is_trans)
-         ctx.trans_since_wr_by_trans.inc();
+         ctx.valu_since_wr_by_trans.inc();
+         if (is_trans)
+            ctx.trans_since_wr_by_trans.inc();
 
-      if (is_trans) {
-         for (Definition& def : instr->definitions) {
-            ctx.valu_since_wr_by_trans.set(def.physReg(), def.bytes());
-            ctx.trans_since_wr_by_trans.set(def.physReg(), def.bytes());
+         if (is_trans) {
+            for (Definition& def : instr->definitions) {
+               ctx.valu_since_wr_by_trans.set(def.physReg(), def.bytes());
+               ctx.trans_since_wr_by_trans.set(def.physReg(), def.bytes());
+            }
          }
-      }
 
-      if (state.program->wave_size == 64) {
-         for (Operand& op : instr->operands) {
-            if (op.isLiteral() || (!op.isConstant() && op.physReg().reg() < 128))
-               ctx.sgpr_read_by_valu_as_lanemask.reset();
-         }
-         switch (instr->opcode) {
-         case aco_opcode::v_addc_co_u32:
-         case aco_opcode::v_subb_co_u32:
-         case aco_opcode::v_subbrev_co_u32:
-         case aco_opcode::v_cndmask_b16:
-         case aco_opcode::v_cndmask_b32:
-         case aco_opcode::v_div_fmas_f32:
-         case aco_opcode::v_div_fmas_f64:
-            if (instr->operands.back().physReg() != exec) {
-               ctx.sgpr_read_by_valu_as_lanemask.set(instr->operands.back().physReg().reg());
-               ctx.sgpr_read_by_valu_as_lanemask.set(instr->operands.back().physReg().reg() + 1);
+         if (state.program->wave_size == 64) {
+            for (Operand& op : instr->operands) {
+               if (op.isLiteral() || (!op.isConstant() && op.physReg().reg() < 128))
+                  ctx.sgpr_read_by_valu_as_lanemask.reset();
+            }
+            switch (instr->opcode) {
+            case aco_opcode::v_addc_co_u32:
+            case aco_opcode::v_subb_co_u32:
+            case aco_opcode::v_subbrev_co_u32:
+            case aco_opcode::v_cndmask_b16:
+            case aco_opcode::v_cndmask_b32:
+            case aco_opcode::v_div_fmas_f32:
+            case aco_opcode::v_div_fmas_f64:
+               if (instr->operands.back().physReg() != exec) {
+                  ctx.sgpr_read_by_valu_as_lanemask.set(instr->operands.back().physReg().reg());
+                  ctx.sgpr_read_by_valu_as_lanemask.set(instr->operands.back().physReg().reg() + 1);
+               }
+               break;
+            default: break;
             }
-            break;
-         default: break;
          }
       }
    }
@@ -1494,14 +1507,23 @@ handle_instruction_gfx11(State& state, NOP_ctx_gfx11& ctx, aco_ptr<Instruction>&
     * Handle LDSDIR writing a VGPR after it's used by a VMEM/DS instruction.
     */
    if (instr->isVMEM() || instr->isFlatLike()) {
-      for (Definition& def : instr->definitions)
-         fill_vgpr_bitset(ctx.vgpr_used_by_vmem_load, def.physReg(), def.bytes());
       if (instr->definitions.empty()) {
          for (Operand& op : instr->operands)
             fill_vgpr_bitset(ctx.vgpr_used_by_vmem_store, op.physReg(), op.bytes());
       } else {
+         uint8_t vmem_type = state.program->gfx_level >= GFX12
+                                ? get_vmem_type(state.program->gfx_level, instr.get())
+                                : vmem_nosampler;
+         std::bitset<256>* vgprs = &ctx.vgpr_used_by_vmem_load;
+         if (vmem_type == vmem_sampler)
+            vgprs = &ctx.vgpr_used_by_vmem_sample;
+         else if (vmem_type == vmem_bvh)
+            vgprs = &ctx.vgpr_used_by_vmem_bvh;
+
+         for (Definition& def : instr->definitions)
+            fill_vgpr_bitset(*vgprs, def.physReg(), def.bytes());
          for (Operand& op : instr->operands)
-            fill_vgpr_bitset(ctx.vgpr_used_by_vmem_load, op.physReg(), op.bytes());
+            fill_vgpr_bitset(*vgprs, op.physReg(), op.bytes());
       }
    }
    if (instr->isDS() || instr->isFlat()) {
@@ -1513,11 +1535,17 @@ handle_instruction_gfx11(State& state, NOP_ctx_gfx11& ctx, aco_ptr<Instruction>&
    wait_imm imm;
    if (instr->isVALU() || instr->isEXP() || vm_vsrc == 0) {
       ctx.vgpr_used_by_vmem_load.reset();
+      ctx.vgpr_used_by_vmem_sample.reset();
+      ctx.vgpr_used_by_vmem_bvh.reset();
       ctx.vgpr_used_by_vmem_store.reset();
       ctx.vgpr_used_by_ds.reset();
    } else if (imm.unpack(state.program->gfx_level, instr.get())) {
       if (imm.vm == 0)
          ctx.vgpr_used_by_vmem_load.reset();
+      if (imm.sample == 0)
+         ctx.vgpr_used_by_vmem_sample.reset();
+      if (imm.bvh == 0)
+         ctx.vgpr_used_by_vmem_bvh.reset();
       if (imm.lgkm == 0)
          ctx.vgpr_used_by_ds.reset();
       if (imm.vs == 0)
@@ -1525,10 +1553,17 @@ handle_instruction_gfx11(State& state, NOP_ctx_gfx11& ctx, aco_ptr<Instruction>&
    }
    if (instr->isLDSDIR()) {
       if (ctx.vgpr_used_by_vmem_load[instr->definitions[0].physReg().reg() - 256] ||
+          ctx.vgpr_used_by_vmem_sample[instr->definitions[0].physReg().reg() - 256] ||
+          ctx.vgpr_used_by_vmem_bvh[instr->definitions[0].physReg().reg() - 256] ||
           ctx.vgpr_used_by_vmem_store[instr->definitions[0].physReg().reg() - 256] ||
           ctx.vgpr_used_by_ds[instr->definitions[0].physReg().reg() - 256]) {
-         bld.sopp(aco_opcode::s_waitcnt_depctr, 0xffe3);
+         if (state.program->gfx_level >= GFX12)
+            instr->ldsdir().wait_vsrc = 0;
+         else
+            bld.sopp(aco_opcode::s_waitcnt_depctr, 0xffe3);
          ctx.vgpr_used_by_vmem_load.reset();
+         ctx.vgpr_used_by_vmem_sample.reset();
+         ctx.vgpr_used_by_vmem_bvh.reset();
          ctx.vgpr_used_by_vmem_store.reset();
          ctx.vgpr_used_by_ds.reset();
       }
@@ -1591,7 +1626,7 @@ resolve_all_gfx11(State& state, NOP_ctx_gfx11& ctx,
    }
 
    /* VALUMaskWriteHazard */
-   if (state.program->wave_size == 64 &&
+   if (state.program->gfx_level < GFX12 && state.program->wave_size == 64 &&
        (ctx.sgpr_read_by_valu_as_lanemask.any() ||
         ctx.sgpr_read_by_valu_as_lanemask_then_wr_by_salu.any())) {
       waitcnt_depctr &= 0xfffe;
@@ -1601,7 +1636,8 @@ resolve_all_gfx11(State& state, NOP_ctx_gfx11& ctx,
 
    /* LdsDirectVMEMHazard */
    if (ctx.vgpr_used_by_vmem_load.any() || ctx.vgpr_used_by_vmem_store.any() ||
-       ctx.vgpr_used_by_ds.any()) {
+       ctx.vgpr_used_by_ds.any() || ctx.vgpr_used_by_vmem_sample.any() ||
+       ctx.vgpr_used_by_vmem_bvh.any()) {
       waitcnt_depctr &= 0xffe3;
       ctx.vgpr_used_by_vmem_load.reset();
       ctx.vgpr_used_by_vmem_store.reset();

src/amd/compiler/aco_instruction_selection.cpp

@@ -3889,14 +3889,6 @@ visit_alu_instr(isel_context* ctx, nir_alu_instr* instr)
    case nir_op_fddy_fine:
    case nir_op_fddx_coarse:
    case nir_op_fddy_coarse: {
-      if (!nir_src_is_divergent(instr->src[0].src)) {
-         /* Source is the same in all lanes, so the derivative is zero.
-          * This also avoids emitting invalid IR.
-          */
-         bld.copy(Definition(dst), Operand::zero(dst.bytes()));
-         break;
-      }
-
       uint16_t dpp_ctrl1, dpp_ctrl2;
       if (instr->op == nir_op_fddx_fine) {
          dpp_ctrl1 = dpp_quad_perm(0, 0, 2, 2);
@@ -3923,8 +3915,12 @@ visit_alu_instr(isel_context* ctx, nir_alu_instr* instr)
          opsel_lo |= opsel_lo << 1;
          opsel_hi |= opsel_hi << 1;
-         Temp tl = bld.vop1_dpp(aco_opcode::v_mov_b32, bld.def(v1), src, dpp_ctrl1);
-         Temp tr = bld.vop1_dpp(aco_opcode::v_mov_b32, bld.def(v1), src, dpp_ctrl2);
+         Temp tl = src;
+         Temp tr = src;
+         if (nir_src_is_divergent(instr->src[0].src)) {
+            tl = bld.vop1_dpp(aco_opcode::v_mov_b32, bld.def(v1), src, dpp_ctrl1);
+            tr = bld.vop1_dpp(aco_opcode::v_mov_b32, bld.def(v1), src, dpp_ctrl2);
+         }
          VALU_instruction& sub =
             bld.vop3p(aco_opcode::v_pk_add_f16, Definition(dst), tr, tl, opsel_lo, opsel_hi)
@@ -3935,9 +3931,11 @@ visit_alu_instr(isel_context* ctx, nir_alu_instr* instr)
       } else {
          Temp src = as_vgpr(ctx, get_alu_src(ctx, instr->src[0]));
-         if (ctx->program->gfx_level >= GFX8) {
-            aco_opcode sub =
-               instr->def.bit_size == 16 ? aco_opcode::v_sub_f16 : aco_opcode::v_sub_f32;
+         aco_opcode sub =
+            instr->def.bit_size == 16 ? aco_opcode::v_sub_f16 : aco_opcode::v_sub_f32;
+         if (!nir_src_is_divergent(instr->src[0].src)) {
+            bld.vop2(sub, Definition(dst), src, src);
+         } else if (ctx->program->gfx_level >= GFX8) {
             Temp tl = bld.vop1_dpp(aco_opcode::v_mov_b32, bld.def(v1), src, dpp_ctrl1);
             bld.vop2_dpp(sub, Definition(dst), src, tl, dpp_ctrl2);
          } else {
@@ -6068,7 +6066,11 @@ static MIMG_instruction*
 emit_mimg(Builder& bld, aco_opcode op, Temp dst, Temp rsrc, Operand samp, std::vector<Temp> coords,
           Operand vdata = Operand(v1))
 {
+   bool is_vsample = !samp.isUndefined() || op == aco_opcode::image_msaa_load;
+
    size_t nsa_size = bld.program->dev.max_nsa_vgprs;
+   if (!is_vsample && bld.program->gfx_level >= GFX12)
+      nsa_size++; /* VIMAGE can encode one more VADDR */
    nsa_size = bld.program->gfx_level >= GFX11 || coords.size() <= nsa_size ? nsa_size : 0;
    const bool strict_wqm = coords[0].regClass().is_linear_vgpr();
@@ -8295,18 +8297,6 @@ visit_intrinsic(isel_context* ctx, nir_intrinsic_instr* instr)
       emit_split_vector(ctx, dst, 3);
       break;
    }
-   case nir_intrinsic_load_ray_launch_size: {
-      Temp dst = get_ssa_temp(ctx, &instr->def);
-      bld.copy(Definition(dst), Operand(get_arg(ctx, ctx->args->rt.launch_size)));
-      emit_split_vector(ctx, dst, 3);
-      break;
-   }
-   case nir_intrinsic_load_ray_launch_id: {
-      Temp dst = get_ssa_temp(ctx, &instr->def);
-      bld.copy(Definition(dst), Operand(get_arg(ctx, ctx->args->rt.launch_id)));
-      emit_split_vector(ctx, dst, 3);
-      break;
-   }
    case nir_intrinsic_load_local_invocation_id: {
       Temp dst = get_ssa_temp(ctx, &instr->def);
       if (ctx->options->gfx_level >= GFX11) {
@@ -8341,11 +8331,8 @@ visit_intrinsic(isel_context* ctx, nir_intrinsic_instr* instr)
    case nir_intrinsic_load_workgroup_id: {
       Temp dst = get_ssa_temp(ctx, &instr->def);
       if (ctx->stage.hw == AC_HW_COMPUTE_SHADER) {
-         const struct ac_arg* ids = ctx->args->workgroup_ids;
-         bld.pseudo(aco_opcode::p_create_vector, Definition(dst),
-                    ids[0].used ? Operand(get_arg(ctx, ids[0])) : Operand::zero(),
-                    ids[1].used ? Operand(get_arg(ctx, ids[1])) : Operand::zero(),
-                    ids[2].used ? Operand(get_arg(ctx, ids[2])) : Operand::zero());
+         bld.pseudo(aco_opcode::p_create_vector, Definition(dst), ctx->workgroup_id[0],
+                    ctx->workgroup_id[1], ctx->workgroup_id[2]);
          emit_split_vector(ctx, dst, 3);
       } else {
          isel_err(&instr->instr, "Unsupported stage for load_workgroup_id");
@@ -8977,6 +8964,15 @@ visit_intrinsic(isel_context* ctx, nir_intrinsic_instr* instr)
    case nir_intrinsic_shader_clock: {
       Temp dst = get_ssa_temp(ctx, &instr->def);
       if (nir_intrinsic_memory_scope(instr) == SCOPE_SUBGROUP &&
+          ctx->options->gfx_level >= GFX12) {
+         Temp hi0 = bld.tmp(s1);
+         Temp hi1 = bld.tmp(s1);
+         Temp lo = bld.tmp(s1);
+         bld.pseudo(aco_opcode::p_shader_cycles_hi_lo_hi, Definition(hi0), Definition(lo), Definition(hi1));
+         Temp hi_eq = bld.sopc(aco_opcode::s_cmp_eq_u32, bld.def(s1, scc), hi0, hi1);
+         lo = bld.sop2(aco_opcode::s_cselect_b32, bld.def(s1), lo, Operand::zero(), bld.scc(hi_eq));
+         bld.pseudo(aco_opcode::p_create_vector, Definition(dst), lo, hi1);
+      } else if (nir_intrinsic_memory_scope(instr) == SCOPE_SUBGROUP &&
           ctx->options->gfx_level >= GFX10_3) {
          /* "((size - 1) << 11) | register" (SHADER_CYCLES is encoded as register 29) */
          Temp clock = bld.sopk(aco_opcode::s_getreg_b32, bld.def(s1), ((20 - 1) << 11) | 29);
@@ -9105,10 +9101,6 @@ visit_intrinsic(isel_context* ctx, nir_intrinsic_instr* instr)
       break;
    }
    case nir_intrinsic_bvh64_intersect_ray_amd: visit_bvh64_intersect_ray_amd(ctx, instr); break;
-   case nir_intrinsic_load_rt_dynamic_callable_stack_base_amd:
-      bld.copy(Definition(get_ssa_temp(ctx, &instr->def)),
-               get_arg(ctx, ctx->args->rt.dynamic_callable_stack_base));
-      break;
    case nir_intrinsic_load_resume_shader_address_amd: {
       bld.pseudo(aco_opcode::p_resume_shader_address, Definition(get_ssa_temp(ctx, &instr->def)),
                  bld.def(s1, scc), Operand::c32(nir_intrinsic_call_idx(instr)));
@@ -11139,6 +11131,9 @@ add_startpgm(struct isel_context* ctx)
          def_count++;
    }
+   if (ctx->stage.hw == AC_HW_COMPUTE_SHADER && ctx->program->gfx_level >= GFX12)
+      def_count += 2;
+
    Instruction* startpgm = create_instruction(aco_opcode::p_startpgm, Format::PSEUDO, 0, def_count);
    ctx->block->instructions.emplace_back(startpgm);
    for (unsigned i = 0, arg = 0; i < ctx->args->arg_count; i++) {
@@ -11171,6 +11166,32 @@ add_startpgm(struct isel_context* ctx)
       }
    }
+   if (ctx->program->gfx_level >= GFX12 && ctx->stage.hw == AC_HW_COMPUTE_SHADER) {
+      Temp idx = ctx->program->allocateTmp(s1);
+      Temp idy = ctx->program->allocateTmp(s1);
+      startpgm->definitions[def_count - 2] = Definition(idx);
+      startpgm->definitions[def_count - 2].setFixed(PhysReg(108 + 9 /*ttmp9*/));
+      startpgm->definitions[def_count - 1] = Definition(idy);
+      startpgm->definitions[def_count - 1].setFixed(PhysReg(108 + 7 /*ttmp7*/));
+      ctx->workgroup_id[0] = Operand(idx);
+      if (ctx->args->workgroup_ids[2].used) {
+         Builder bld(ctx->program, ctx->block);
+         ctx->workgroup_id[1] =
+            bld.pseudo(aco_opcode::p_extract, bld.def(s1), bld.def(s1, scc), idy, Operand::zero(),
+                       Operand::c32(16u), Operand::zero());
+         ctx->workgroup_id[2] =
+            bld.pseudo(aco_opcode::p_extract, bld.def(s1), bld.def(s1, scc), idy, Operand::c32(1u),
+                       Operand::c32(16u), Operand::zero());
+      } else {
+         ctx->workgroup_id[1] = Operand(idy);
+         ctx->workgroup_id[2] = Operand::zero();
+      }
+   } else if (ctx->stage.hw == AC_HW_COMPUTE_SHADER) {
+      const struct ac_arg* ids = ctx->args->workgroup_ids;
+      for (unsigned i = 0; i < 3; i++)
+         ctx->workgroup_id[i] = ids[i].used ? Operand(get_arg(ctx, ids[i])) : Operand::zero();
+   }
+
    /* epilog has no scratch */
    if (ctx->args->scratch_offset.used) {
       if (ctx->program->gfx_level < GFX9) {
@@ -12280,10 +12301,18 @@ select_rt_prolog(Program* program, ac_shader_config* config,
    PhysReg in_sbt_desc = get_arg_reg(in_args, in_args->rt.sbt_descriptors);
    PhysReg in_launch_size_addr = get_arg_reg(in_args, in_args->rt.launch_size_addr);
    PhysReg in_stack_base = get_arg_reg(in_args, in_args->rt.dynamic_callable_stack_base);
-   PhysReg in_wg_id_x = get_arg_reg(in_args, in_args->workgroup_ids[0]);
-   PhysReg in_wg_id_y = get_arg_reg(in_args, in_args->workgroup_ids[1]);
-   PhysReg in_wg_id_z = get_arg_reg(in_args, in_args->workgroup_ids[2]);
+   PhysReg in_wg_id_x;
+   PhysReg in_wg_id_y;
+   PhysReg in_wg_id_z;
    PhysReg in_scratch_offset;
+   if (options->gfx_level < GFX12) {
+      in_wg_id_x = get_arg_reg(in_args, in_args->workgroup_ids[0]);
+      in_wg_id_y = get_arg_reg(in_args, in_args->workgroup_ids[1]);
+      in_wg_id_z = get_arg_reg(in_args, in_args->workgroup_ids[2]);
+   } else {
+      in_wg_id_x = PhysReg(108 + 9 /*ttmp9*/);
+      in_wg_id_y = PhysReg(108 + 7 /*ttmp7*/);
+   }
    if (options->gfx_level < GFX11)
       in_scratch_offset = get_arg_reg(in_args, in_args->scratch_offset);
    PhysReg in_local_ids[2] = {
@@ -12306,12 +12335,12 @@ select_rt_prolog(Program* program, ac_shader_config* config,
     * Shader Record Ptr:           v[6-7]
     */
    PhysReg out_uniform_shader_addr = get_arg_reg(out_args, out_args->rt.uniform_shader_addr);
-   PhysReg out_launch_size_x = get_arg_reg(out_args, out_args->rt.launch_size);
-   PhysReg out_launch_size_y = out_launch_size_x.advance(4);
-   PhysReg out_launch_size_z = out_launch_size_y.advance(4);
+   PhysReg out_launch_size_x = get_arg_reg(out_args, out_args->rt.launch_sizes[0]);
+   PhysReg out_launch_size_y = get_arg_reg(out_args, out_args->rt.launch_sizes[1]);
+   PhysReg out_launch_size_z = get_arg_reg(out_args, out_args->rt.launch_sizes[2]);
    PhysReg out_launch_ids[3];
    for (unsigned i = 0; i < 3; i++)
-      out_launch_ids[i] = get_arg_reg(out_args, out_args->rt.launch_id).advance(i * 4);
+      out_launch_ids[i] = get_arg_reg(out_args, out_args->rt.launch_ids[i]);
    PhysReg out_stack_ptr = get_arg_reg(out_args, out_args->rt.dynamic_callable_stack_base);
    PhysReg out_record_ptr = get_arg_reg(out_args, out_args->rt.shader_record);
@@ -12321,6 +12350,8 @@ select_rt_prolog(Program* program, ac_shader_config* config,
    num_sgprs += 2;
    PhysReg tmp_ring_offsets = PhysReg{num_sgprs};
    num_sgprs += 2;
+   PhysReg tmp_wg_id_x_times_size = PhysReg{num_sgprs};
+   num_sgprs++;
    PhysReg tmp_invocation_idx = PhysReg{256 + num_vgprs++};
@@ -12370,9 +12401,18 @@ select_rt_prolog(Program* program, ac_shader_config* config,
                Operand(in_local_ids[0], v1));
    }
    /* Do this backwards to reduce some RAW hazards on GFX11+ */
-   bld.vop1(aco_opcode::v_mov_b32, Definition(out_launch_ids[2], v1), Operand(in_wg_id_z, s1));
-   bld.vop3(aco_opcode::v_mad_u32_u24, Definition(out_launch_ids[1], v1), Operand(in_wg_id_y, s1),
-            Operand::c32(program->workgroup_size == 32 ? 4 : 8), Operand(in_local_ids[1], v1));
+   if (options->gfx_level >= GFX12) {
+      bld.vop2_e64(aco_opcode::v_lshrrev_b32, Definition(out_launch_ids[2], v1), Operand::c32(16),
+                   Operand(in_wg_id_y, s1));
+      bld.vop3(aco_opcode::v_mad_u32_u16, Definition(out_launch_ids[1], v1),
+               Operand(in_wg_id_y, s1), Operand::c32(program->workgroup_size == 32 ? 4 : 8),
+               Operand(in_local_ids[1], v1));
+   } else {
+      bld.vop1(aco_opcode::v_mov_b32, Definition(out_launch_ids[2], v1), Operand(in_wg_id_z, s1));
+      bld.vop3(aco_opcode::v_mad_u32_u24, Definition(out_launch_ids[1], v1),
+               Operand(in_wg_id_y, s1), Operand::c32(program->workgroup_size == 32 ? 4 : 8),
+               Operand(in_local_ids[1], v1));
+   }
    bld.vop3(aco_opcode::v_mad_u32_u24, Definition(out_launch_ids[0], v1), Operand(in_wg_id_x, s1),
             Operand::c32(8), Operand(in_local_ids[0], v1));
@@ -12398,14 +12438,14 @@ select_rt_prolog(Program* program, ac_shader_config* config,
    /* For 1D dispatches converted into 2D ones, we need to fix up the launch IDs.
     * Calculating the 1D launch ID is: id = local_invocation_index + (wg_id.x * wg_size).
-    * in_wg_id_x now holds wg_id.x * wg_size.
+    * tmp_wg_id_x_times_size now holds wg_id.x * wg_size.
     */
-   bld.sop2(aco_opcode::s_lshl_b32, Definition(in_wg_id_x, s1), Definition(scc, s1),
+   bld.sop2(aco_opcode::s_lshl_b32, Definition(tmp_wg_id_x_times_size, s1), Definition(scc, s1),
             Operand(in_wg_id_x, s1), Operand::c32(program->workgroup_size == 32 ? 5 : 6));
    /* Calculate and add local_invocation_index */
    bld.vop3(aco_opcode::v_mbcnt_lo_u32_b32, Definition(tmp_invocation_idx, v1), Operand::c32(-1u),
-            Operand(in_wg_id_x, s1));
+            Operand(tmp_wg_id_x_times_size, s1));
    if (program->wave_size == 64) {
       if (program->gfx_level <= GFX7)
          bld.vop2(aco_opcode::v_mbcnt_hi_u32_b32, Definition(tmp_invocation_idx, v1),

src/amd/compiler/aco_instruction_selection.h

@@ -72,6 +72,7 @@ struct isel_context {
    nir_unsigned_upper_bound_config ub_config;
 
    Temp arg_temps[AC_MAX_ARGS];
+   Operand workgroup_id[3];
 
    /* tessellation information */
    uint64_t tcs_temp_only_inputs;

src/amd/compiler/aco_instruction_selection_setup.cpp

@@ -442,7 +442,6 @@ init_context(isel_context* ctx, nir_shader* shader)
                case nir_intrinsic_load_push_constant:
                case nir_intrinsic_load_workgroup_id:
                case nir_intrinsic_load_num_workgroups:
-               case nir_intrinsic_load_ray_launch_size:
                case nir_intrinsic_load_sbt_base_amd:
                case nir_intrinsic_load_subgroup_id:
                case nir_intrinsic_load_num_subgroups:
@@ -481,7 +480,6 @@ init_context(isel_context* ctx, nir_shader* shader)
                case nir_intrinsic_load_local_invocation_id:
                case nir_intrinsic_load_local_invocation_index:
                case nir_intrinsic_load_subgroup_invocation:
-               case nir_intrinsic_load_ray_launch_id:
                case nir_intrinsic_load_tess_coord:
                case nir_intrinsic_write_invocation_amd:
                case nir_intrinsic_mbcnt_amd:
@@ -505,7 +503,6 @@ init_context(isel_context* ctx, nir_shader* shader)
                case nir_intrinsic_gds_atomic_add_amd:
                case nir_intrinsic_bvh64_intersect_ray_amd:
                case nir_intrinsic_load_vector_arg_amd:
-               case nir_intrinsic_load_rt_dynamic_callable_stack_base_amd:
                case nir_intrinsic_ordered_xfb_counter_add_gfx11_amd:
                case nir_intrinsic_cmat_muladd_amd:
                case nir_intrinsic_unit_test_divergent_amd: type = RegType::vgpr; break;

src/amd/compiler/aco_ir.cpp

@@ -155,7 +155,8 @@ init_program(Program* program, Stage stage, const struct aco_shader_info* info,
    if (program->family == CHIP_TAHITI || program->family == CHIP_CARRIZO ||
        program->family == CHIP_HAWAII)
       program->dev.has_fast_fma32 = true;
-   program->dev.has_mac_legacy32 = program->gfx_level <= GFX7 || program->gfx_level >= GFX10;
+   program->dev.has_mac_legacy32 = program->gfx_level <= GFX7 || program->gfx_level == GFX10;
+   program->dev.has_fmac_legacy32 = program->gfx_level >= GFX10_3 && program->gfx_level < GFX12;
 
    program->dev.fused_mad_mix = program->gfx_level >= GFX10;
    if (program->family == CHIP_VEGA12 || program->family == CHIP_VEGA20 ||
@@ -174,7 +175,10 @@ init_program(Program* program, Stage stage, const struct aco_shader_info* info,
       program->dev.scratch_global_offset_max = 4095;
    }
 
-   if (program->gfx_level >= GFX11) {
+   if (program->gfx_level >= GFX12) {
+      /* Same as GFX11, except one less for VSAMPLE. */
+      program->dev.max_nsa_vgprs = 3;
+   } else if (program->gfx_level >= GFX11) {
       /* GFX11 can have only 1 NSA dword. The last VGPR isn't included here because it contains the
        * rest of the address.
        */

src/amd/compiler/aco_ir.h

@@ -2011,6 +2011,7 @@ struct DeviceInfo {
    unsigned simd_per_cu;
    bool has_fast_fma32 = false;
    bool has_mac_legacy32 = false;
+   bool has_fmac_legacy32 = false;
    bool fused_mad_mix = false;
    bool xnack_enabled = false;
    bool sram_ecc_enabled = false;

src/amd/compiler/aco_opcodes.py

@@ -445,6 +445,8 @@ insn("p_dual_src_export_gfx11")
 # shader to pass arguments to next part.
 insn("p_end_with_regs")
 
+insn("p_shader_cycles_hi_lo_hi")
+
 # SOP2 instructions: 2 scalar inputs, 1 scalar output (+optional scc)
 SOP2 = {
    ("s_add_u32",            dst(1, SCC), src(1, 1), op(0x00)),

src/amd/compiler/aco_optimizer.cpp

@@ -2355,11 +2355,11 @@ optimize_cmp_subgroup_invocation(opt_ctx& ctx, aco_ptr<Instruction>& instr)
 
    Instruction* cpy = NULL;
    const uint64_t mask = BITFIELD64_RANGE(first_bit, num_bits);
-   if (wave_size == 64 && mask > 0x7fffffff && mask != -1ull) {
-      /* Mask can't be represented as a 64-bit constant or literal, use s_bfm_b64. */
-      cpy = create_instruction(aco_opcode::s_bfm_b64, Format::SOP2, 2, 1);
-      cpy->operands[0] = Operand::c32(num_bits);
-      cpy->operands[1] = Operand::c32(first_bit);
+   if (!Operand::is_constant_representable(mask, wave_size / 8, true, false)) {
+      /* Mask can't be represented as a 64-bit constant or literal, create a vector */
+      cpy = create_instruction(aco_opcode::p_create_vector, Format::PSEUDO, 2, 1);
+      cpy->operands[0] = Operand::c32(mask);
+      cpy->operands[1] = Operand::c32(mask >> 32);
    } else {
       /* Copy mask as a literal constant. */
       cpy = create_instruction(aco_opcode::p_parallelcopy, Format::PSEUDO, 1, 1);

src/amd/compiler/aco_register_allocation.cpp

@@ -2089,6 +2089,8 @@ operand_can_use_reg(amd_gfx_level gfx_level, aco_ptr<Instruction>& instr, unsign
              (reg != m0 || idx == 1 || idx == 3) && /* offset can be m0 */
              (reg != vcc || (instr->definitions.empty() && idx == 2) ||
               gfx_level >= GFX10); /* sdata can be vcc */
+   case Format::MUBUF:
+   case Format::MTBUF: return idx != 2 || gfx_level < GFX12 || reg != scc;
    default:
       // TODO: there are more instructions with restrictions on registers
       return true;
@@ -2573,7 +2575,7 @@ get_affinities(ra_ctx& ctx, std::vector<IDSet>& live_out_per_block)
                   ctx.vectors[op.tempId()] = instr.get();
             }
          } else if (instr->format == Format::MIMG && instr->operands.size() > 4 &&
-                    !instr->mimg().strict_wqm) {
+                    !instr->mimg().strict_wqm && ctx.program->gfx_level < GFX12) {
             for (unsigned i = 3; i < instr->operands.size(); i++)
                ctx.vectors[instr->operands[i].tempId()] = instr.get();
          } else if (instr->opcode == aco_opcode::p_split_vector &&
@@ -2641,11 +2643,15 @@ get_affinities(ra_ctx& ctx, std::vector<IDSet>& live_out_per_block)
                   break;
 
                case aco_opcode::v_mad_legacy_f32:
-               case aco_opcode::v_fma_legacy_f32:
                   if (instr->usesModifiers() || !ctx.program->dev.has_mac_legacy32)
                      continue;
                   op = instr->operands[2];
                   break;
+               case aco_opcode::v_fma_legacy_f32:
+                  if (instr->usesModifiers() || !ctx.program->dev.has_fmac_legacy32)
+                     continue;
+                  op = instr->operands[2];
+                  break;
 
                default: continue;
                }
@@ -2741,7 +2747,7 @@ optimize_encoding_vop2(Program* program, ra_ctx& ctx, RegisterFile& register_fil
         (instr->opcode != aco_opcode::v_fma_f16 || program->gfx_level < GFX10) &&
         (instr->opcode != aco_opcode::v_pk_fma_f16 || program->gfx_level < GFX10) &&
         (instr->opcode != aco_opcode::v_mad_legacy_f32 || !program->dev.has_mac_legacy32) &&
-        (instr->opcode != aco_opcode::v_fma_legacy_f32 || !program->dev.has_mac_legacy32) &&
+        (instr->opcode != aco_opcode::v_fma_legacy_f32 || !program->dev.has_fmac_legacy32) &&
         (instr->opcode != aco_opcode::v_dot4_i32_i8 || program->family == CHIP_VEGA20)) ||
        !instr->operands[2].isTemp() || !instr->operands[2].isKillBeforeDef() ||
        instr->operands[2].getTemp().type() != RegType::vgpr ||

src/amd/compiler/aco_scheduler.cpp

@@ -596,6 +596,7 @@ perform_hazard_query(hazard_query* query, Instruction* instr, bool upwards)
    /* don't move non-reorderable instructions */
    if (instr->opcode == aco_opcode::s_memtime || instr->opcode == aco_opcode::s_memrealtime ||
        instr->opcode == aco_opcode::s_setprio || instr->opcode == aco_opcode::s_getreg_b32 ||
+       instr->opcode == aco_opcode::p_shader_cycles_hi_lo_hi ||
        instr->opcode == aco_opcode::p_init_scratch ||
        instr->opcode == aco_opcode::p_jump_to_epilog ||
        instr->opcode == aco_opcode::s_sendmsg_rtn_b32 ||

src/amd/compiler/aco_scheduler_ilp.cpp

@@ -97,6 +97,12 @@ can_reorder(const Instruction* const instr)
    case aco_opcode::s_set_gpr_idx_idx:
    case aco_opcode::s_sendmsg_rtn_b32:
    case aco_opcode::s_sendmsg_rtn_b64:
+   case aco_opcode::s_barrier_signal:
+   case aco_opcode::s_barrier_signal_isfirst:
+   case aco_opcode::s_get_barrier_state:
+   case aco_opcode::s_barrier_init:
+   case aco_opcode::s_barrier_join:
+   case aco_opcode::s_wakeup_barrier:
    /* SOPK */
    case aco_opcode::s_cbranch_i_fork:
    case aco_opcode::s_getreg_b32:

src/amd/compiler/aco_validate.cpp

@@ -722,6 +722,9 @@ validate_ir(Program* program)
                   instr.get());
             check(instr->operands.size() < 4 || instr->operands[3].isOfType(RegType::vgpr),
                   "VMEM write data must be vgpr", instr.get());
+            if (instr->operands.size() >= 3 && instr->operands[2].isConstant())
+               check(program->gfx_level < GFX12 || instr->operands[2].constantValue() == 0,
+                     "VMEM SOFFSET must not be non-zero constant on GFX12+", instr.get());
 
             const bool d16 =
                instr->opcode ==
@@ -800,8 +803,11 @@ validate_ir(Program* program)
                         check(instr->operands[i].regClass() == v1,
                               "GFX10 MIMG VADDR must be v1 if NSA is used", instr.get());
                      } else {
+                        unsigned num_scalar =
+                           program->gfx_level >= GFX12 ? (instr->operands.size() - 4) : 4;
                         if (instr->opcode != aco_opcode::image_bvh_intersect_ray &&
-                            instr->opcode != aco_opcode::image_bvh64_intersect_ray && i < 7) {
+                            instr->opcode != aco_opcode::image_bvh64_intersect_ray &&
+                            i < 3 + num_scalar) {
                            check(instr->operands[i].regClass() == v1,
                                  "first 4 GFX11 MIMG VADDR must be v1 if NSA is used", instr.get());
                         }

src/amd/compiler/tests/test_assembler.cpp

@@ -85,35 +85,39 @@ BEGIN_TEST(assembler.long_jump.unconditional_forwards)
 END_TEST
 
 BEGIN_TEST(assembler.long_jump.conditional_forwards)
-   if (!setup_cs(NULL, (amd_gfx_level)GFX10))
-      return;
-
-   //! BB0:
-   //! s_cbranch_scc1 BB1                                          ; bf850006
-   //! s_getpc_b64 s[0:1]                                          ; be801f00
-   //! s_addc_u32 s0, s0, 0x20014                                  ; 8200ff00 00020014
-   //! s_bitcmp1_b32 s0, 0                                         ; bf0d8000
-   //! s_bitset0_b32 s0, 0                                         ; be801b80
-   //! s_setpc_b64 s[0:1]                                          ; be802000
-   bld.sopp(aco_opcode::s_cbranch_scc0, Definition(PhysReg(0), s2), 2);
-
-   bld.reset(program->create_and_insert_block());
-
-   //! BB1:
-   //! s_nop 0 ; bf800000
-   //!(then repeated 32767 times)
-   for (unsigned i = 0; i < INT16_MAX + 1; i++)
-      bld.sopp(aco_opcode::s_nop, 0);
-
-   //! BB2:
-   //! s_endpgm                                                    ; bf810000
-   bld.reset(program->create_and_insert_block());
+   for (amd_gfx_level gfx : filter_gfx_levels({GFX10, GFX12})) {
+      if (!setup_cs(NULL, gfx))
+         continue;
 
-   program->blocks[1].linear_preds.push_back(0u);
-   program->blocks[2].linear_preds.push_back(0u);
-   program->blocks[2].linear_preds.push_back(1u);
+      //! BB0:
+      //! s_cbranch_scc1 BB1                                          ; $_
+      //! s_getpc_b64 s[0:1]                                          ; $_
+      //~gfx12! s_sext_i32_i16 s1, s1                                 ; $_
+      //~gfx10! s_addc_u32 s0, s0, 0x20014                            ; $_ $_
+      //~gfx12! s_add_co_ci_u32 s0, s0, 0x20014                       ; $_ $_
+      //! s_bitcmp1_b32 s0, 0                                         ; $_
+      //! s_bitset0_b32 s0, 0                                         ; $_
+      //! s_setpc_b64 s[0:1]                                          ; $_
+      bld.sopp(aco_opcode::s_cbranch_scc0, Definition(PhysReg(0), s2), 2);
+
+      bld.reset(program->create_and_insert_block());
+
+      //! BB1:
+      //! s_nop 0 ; bf800000
+      //!(then repeated 32767 times)
+      for (unsigned i = 0; i < INT16_MAX + 1; i++)
+         bld.sopp(aco_opcode::s_nop, 0);
+
+      //! BB2:
+      //! s_endpgm                                                    ; $_
+      bld.reset(program->create_and_insert_block());
+
+      program->blocks[1].linear_preds.push_back(0u);
+      program->blocks[2].linear_preds.push_back(0u);
+      program->blocks[2].linear_preds.push_back(1u);
 
-   finish_assembler_test();
+      finish_assembler_test();
+   }
 END_TEST
 
 BEGIN_TEST(assembler.long_jump.unconditional_backwards)

src/amd/llvm/ac_nir_to_llvm.c

@@ -3827,7 +3827,7 @@ static bool visit_intrinsic(struct ac_nir_context *ctx, nir_intrinsic_instr *ins
       break;
    }
    case nir_intrinsic_nop_amd: {
-      LLVMValueRef arg = LLVMConstInt(ctx->ac.i32, nir_intrinsic_base(instr), 0);
+      LLVMValueRef arg = LLVMConstInt(ctx->ac.i16, nir_intrinsic_base(instr), 0);
       ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.s.nop", ctx->ac.voidt, &arg, 1, 0);
       break;
    }

src/amd/vulkan/radv_cmd_buffer.h

@@ -78,64 +78,6 @@ enum radv_dynamic_state_bits {
    RADV_DYNAMIC_ALL = (1ull << 52) - 1,
 };
 
-enum radv_cmd_dirty_dynamic_bits {
-   /* Keep the dynamic state dirty bits in sync with
-    * enum radv_dynamic_state_bits */
-   RADV_CMD_DIRTY_DYNAMIC_VIEWPORT = 1ull << 0,
-   RADV_CMD_DIRTY_DYNAMIC_SCISSOR = 1ull << 1,
-   RADV_CMD_DIRTY_DYNAMIC_LINE_WIDTH = 1ull << 2,
-   RADV_CMD_DIRTY_DYNAMIC_DEPTH_BIAS = 1ull << 3,
-   RADV_CMD_DIRTY_DYNAMIC_BLEND_CONSTANTS = 1ull << 4,
-   RADV_CMD_DIRTY_DYNAMIC_DEPTH_BOUNDS = 1ull << 5,
-   RADV_CMD_DIRTY_DYNAMIC_STENCIL_COMPARE_MASK = 1ull << 6,
-   RADV_CMD_DIRTY_DYNAMIC_STENCIL_WRITE_MASK = 1ull << 7,
-   RADV_CMD_DIRTY_DYNAMIC_STENCIL_REFERENCE = 1ull << 8,
-   RADV_CMD_DIRTY_DYNAMIC_DISCARD_RECTANGLE = 1ull << 9,
-   RADV_CMD_DIRTY_DYNAMIC_SAMPLE_LOCATIONS = 1ull << 10,
-   RADV_CMD_DIRTY_DYNAMIC_LINE_STIPPLE = 1ull << 11,
-   RADV_CMD_DIRTY_DYNAMIC_CULL_MODE = 1ull << 12,
-   RADV_CMD_DIRTY_DYNAMIC_FRONT_FACE = 1ull << 13,
-   RADV_CMD_DIRTY_DYNAMIC_PRIMITIVE_TOPOLOGY = 1ull << 14,
-   RADV_CMD_DIRTY_DYNAMIC_DEPTH_TEST_ENABLE = 1ull << 15,
-   RADV_CMD_DIRTY_DYNAMIC_DEPTH_WRITE_ENABLE = 1ull << 16,
-   RADV_CMD_DIRTY_DYNAMIC_DEPTH_COMPARE_OP = 1ull << 17,
-   RADV_CMD_DIRTY_DYNAMIC_DEPTH_BOUNDS_TEST_ENABLE = 1ull << 18,
-   RADV_CMD_DIRTY_DYNAMIC_STENCIL_TEST_ENABLE = 1ull << 19,
-   RADV_CMD_DIRTY_DYNAMIC_STENCIL_OP = 1ull << 20,
-   RADV_CMD_DIRTY_DYNAMIC_VERTEX_INPUT_BINDING_STRIDE = 1ull << 21,
-   RADV_CMD_DIRTY_DYNAMIC_FRAGMENT_SHADING_RATE = 1ull << 22,
-   RADV_CMD_DIRTY_DYNAMIC_PATCH_CONTROL_POINTS = 1ull << 23,
-   RADV_CMD_DIRTY_DYNAMIC_RASTERIZER_DISCARD_ENABLE = 1ull << 24,
-   RADV_CMD_DIRTY_DYNAMIC_DEPTH_BIAS_ENABLE = 1ull << 25,
-   RADV_CMD_DIRTY_DYNAMIC_LOGIC_OP = 1ull << 26,
-   RADV_CMD_DIRTY_DYNAMIC_PRIMITIVE_RESTART_ENABLE = 1ull << 27,
-   RADV_CMD_DIRTY_DYNAMIC_COLOR_WRITE_ENABLE = 1ull << 28,
-   RADV_CMD_DIRTY_DYNAMIC_VERTEX_INPUT = 1ull << 29,
-   RADV_CMD_DIRTY_DYNAMIC_POLYGON_MODE = 1ull << 30,
-   RADV_CMD_DIRTY_DYNAMIC_TESS_DOMAIN_ORIGIN = 1ull << 31,
-   RADV_CMD_DIRTY_DYNAMIC_LOGIC_OP_ENABLE = 1ull << 32,
-   RADV_CMD_DIRTY_DYNAMIC_LINE_STIPPLE_ENABLE = 1ull << 33,
-   RADV_CMD_DIRTY_DYNAMIC_ALPHA_TO_COVERAGE_ENABLE = 1ull << 34,
-   RADV_CMD_DIRTY_DYNAMIC_SAMPLE_MASK = 1ull << 35,
-   RADV_CMD_DIRTY_DYNAMIC_DEPTH_CLIP_ENABLE = 1ull << 36,
-   RADV_CMD_DIRTY_DYNAMIC_CONSERVATIVE_RAST_MODE = 1ull << 37,
-   RADV_CMD_DIRTY_DYNAMIC_DEPTH_CLIP_NEGATIVE_ONE_TO_ONE = 1ull << 38,
-   RADV_CMD_DIRTY_DYNAMIC_PROVOKING_VERTEX_MODE = 1ull << 39,
-   RADV_CMD_DIRTY_DYNAMIC_DEPTH_CLAMP_ENABLE = 1ull << 40,
-   RADV_CMD_DIRTY_DYNAMIC_COLOR_WRITE_MASK = 1ull << 41,
-   RADV_CMD_DIRTY_DYNAMIC_COLOR_BLEND_ENABLE = 1ull << 42,
-   RADV_CMD_DIRTY_DYNAMIC_RASTERIZATION_SAMPLES = 1ull << 43,
-   RADV_CMD_DIRTY_DYNAMIC_LINE_RASTERIZATION_MODE = 1ull << 44,
-   RADV_CMD_DIRTY_DYNAMIC_COLOR_BLEND_EQUATION = 1ull << 45,
-   RADV_CMD_DIRTY_DYNAMIC_DISCARD_RECTANGLE_ENABLE = 1ull << 46,
-   RADV_CMD_DIRTY_DYNAMIC_DISCARD_RECTANGLE_MODE = 1ull << 47,
-   RADV_CMD_DIRTY_DYNAMIC_ATTACHMENT_FEEDBACK_LOOP_ENABLE = 1ull << 48,
-   RADV_CMD_DIRTY_DYNAMIC_SAMPLE_LOCATIONS_ENABLE = 1ull << 49,
-   RADV_CMD_DIRTY_DYNAMIC_ALPHA_TO_ONE_ENABLE = 1ull << 50,
-   RADV_CMD_DIRTY_DYNAMIC_COLOR_ATTACHMENT_MAP = 1ull << 51,
-   RADV_CMD_DIRTY_DYNAMIC_ALL = (1ull << 52) - 1,
-};
-
 enum radv_cmd_dirty_bits {
    RADV_CMD_DIRTY_PIPELINE = 1ull << 0,
    RADV_CMD_DIRTY_INDEX_BUFFER = 1ull << 1,