diff --git a/src/panfrost/midgard/compiler.h b/src/panfrost/midgard/compiler.h
index 52b96a59b2d..b0626d4e389 100644
--- a/src/panfrost/midgard/compiler.h
+++ b/src/panfrost/midgard/compiler.h
@@ -516,34 +516,34 @@ void mir_rewrite_index_src_swizzle(compiler_context *ctx, unsigned old,
                                    unsigned new, unsigned *swizzle);
 bool mir_single_use(compiler_context *ctx, unsigned value);
 unsigned mir_use_count(compiler_context *ctx, unsigned value);
-uint16_t mir_bytemask_of_read_components(midgard_instruction *ins,
+uint16_t mir_bytemask_of_read_components(const midgard_instruction *ins,
                                          unsigned node);
-uint16_t mir_bytemask_of_read_components_index(midgard_instruction *ins,
+uint16_t mir_bytemask_of_read_components_index(const midgard_instruction *ins,
                                                unsigned i);
 uint16_t mir_from_bytemask(uint16_t bytemask, unsigned bits);
-uint16_t mir_bytemask(midgard_instruction *ins);
+uint16_t mir_bytemask(const midgard_instruction *ins);
 uint16_t mir_round_bytemask_up(uint16_t mask, unsigned bits);
 void mir_set_bytemask(midgard_instruction *ins, uint16_t bytemask);
-signed mir_upper_override(midgard_instruction *ins, unsigned inst_size);
+signed mir_upper_override(const midgard_instruction *ins, unsigned inst_size);
 unsigned mir_components_for_type(nir_alu_type T);
-unsigned max_bitsize_for_alu(midgard_instruction *ins);
+unsigned max_bitsize_for_alu(const midgard_instruction *ins);
 midgard_reg_mode reg_mode_for_bitsize(unsigned bitsize);
 
 /* MIR printing */
 
-void mir_print_instruction(midgard_instruction *ins);
-void mir_print_bundle(midgard_bundle *ctx);
-void mir_print_block(midgard_block *block);
-void mir_print_shader(compiler_context *ctx);
-bool mir_nontrivial_mod(midgard_instruction *ins, unsigned i,
+void mir_print_instruction(const midgard_instruction *ins);
+void mir_print_bundle(const midgard_bundle *ctx);
+void mir_print_block(const midgard_block *block);
+void mir_print_shader(const compiler_context *ctx);
+bool mir_nontrivial_mod(const midgard_instruction *ins, unsigned i,
                         bool check_swizzle);
-bool mir_nontrivial_outmod(midgard_instruction *ins);
+bool mir_nontrivial_outmod(const midgard_instruction *ins);
 
 midgard_instruction *mir_insert_instruction_before_scheduled(
-   compiler_context *ctx, midgard_block *block, midgard_instruction *tag,
+   compiler_context *ctx, midgard_block *block, const midgard_instruction *tag,
    midgard_instruction ins);
 midgard_instruction *mir_insert_instruction_after_scheduled(
-   compiler_context *ctx, midgard_block *block, midgard_instruction *tag,
+   compiler_context *ctx, midgard_block *block, const midgard_instruction *tag,
    midgard_instruction ins);
 void mir_flip(midgard_instruction *ins);
 void mir_compute_temp_count(compiler_context *ctx);
@@ -635,7 +635,7 @@ v_load_store_scratch(unsigned srcdest, unsigned index, bool is_store,
 }
 
 static inline bool
-mir_has_arg(midgard_instruction *ins, unsigned arg)
+mir_has_arg(const midgard_instruction *ins, unsigned arg)
 {
    if (!ins)
       return false;
@@ -655,12 +655,12 @@ void midgard_schedule_program(compiler_context *ctx);
 void mir_ra(compiler_context *ctx);
 void mir_squeeze_index(compiler_context *ctx);
 void mir_lower_special_reads(compiler_context *ctx);
-void mir_liveness_ins_update(uint16_t *live, midgard_instruction *ins,
+void mir_liveness_ins_update(uint16_t *live, const midgard_instruction *ins,
                              unsigned max);
 void mir_compute_liveness(compiler_context *ctx);
 void mir_invalidate_liveness(compiler_context *ctx);
-bool mir_is_live_after(compiler_context *ctx, midgard_block *block,
-                       midgard_instruction *start, int src);
+bool mir_is_live_after(compiler_context *ctx, const midgard_block *block,
+                       const midgard_instruction *start, int src);
 
 void mir_create_pipeline_registers(compiler_context *ctx);
 void midgard_promote_uniforms(compiler_context *ctx);
diff --git a/src/panfrost/midgard/helpers.h b/src/panfrost/midgard/helpers.h
index 5cc63e705fb..79954b5c012 100644
--- a/src/panfrost/midgard/helpers.h
+++ b/src/panfrost/midgard/helpers.h
@@ -408,7 +408,8 @@ midgard_is_branch_unit(unsigned unit)
 
 /* Packs ALU mod argument */
 struct midgard_instruction;
-unsigned mir_pack_mod(struct midgard_instruction *ins, unsigned i, bool scalar);
+unsigned mir_pack_mod(const struct midgard_instruction *ins, unsigned i,
+                      bool scalar);
 
 void mir_print_constant_component(FILE *fp, const midgard_constants *consts,
                                   unsigned c, midgard_reg_mode reg_mode,
diff --git a/src/panfrost/midgard/midgard_compile.c b/src/panfrost/midgard/midgard_compile.c
index 9f9c93ccd50..0698bf115be 100644
--- a/src/panfrost/midgard/midgard_compile.c
+++ b/src/panfrost/midgard/midgard_compile.c
@@ -2452,7 +2452,7 @@ inline_alu_constants(compiler_context *ctx, midgard_block *block)
 }
 
 unsigned
-max_bitsize_for_alu(midgard_instruction *ins)
+max_bitsize_for_alu(const midgard_instruction *ins)
 {
    unsigned max_bitsize = 0;
    for (int i = 0; i < MIR_SRC_COUNT; i++) {
diff --git a/src/panfrost/midgard/midgard_emit.c b/src/panfrost/midgard/midgard_emit.c
index 7839760ba5e..09d2fe0ec06 100644
--- a/src/panfrost/midgard/midgard_emit.c
+++ b/src/panfrost/midgard/midgard_emit.c
@@ -78,7 +78,7 @@ midgard_unpack_varying_params(midgard_load_store_word word)
 }
 
 unsigned
-mir_pack_mod(midgard_instruction *ins, unsigned i, bool scalar)
+mir_pack_mod(const midgard_instruction *ins, unsigned i, bool scalar)
 {
    bool integer = midgard_is_integer_op(ins->op);
    unsigned base_size = max_bitsize_for_alu(ins);
diff --git a/src/panfrost/midgard/midgard_liveness.c b/src/panfrost/midgard/midgard_liveness.c
index 984c95f1bcd..443d7ddea26 100644
--- a/src/panfrost/midgard/midgard_liveness.c
+++ b/src/panfrost/midgard/midgard_liveness.c
@@ -25,7 +25,8 @@
 #include "compiler.h"
 
 void
-mir_liveness_ins_update(uint16_t *live, midgard_instruction *ins, unsigned max)
+mir_liveness_ins_update(uint16_t *live, const midgard_instruction *ins,
+                        unsigned max)
 {
    /* live_in[s] = GEN[s] + (live_out[s] - KILL[s]) */
 
@@ -76,8 +77,8 @@ mir_invalidate_liveness(compiler_context *ctx)
 }
 
 bool
-mir_is_live_after(compiler_context *ctx, midgard_block *block,
-                  midgard_instruction *start, int src)
+mir_is_live_after(compiler_context *ctx, const midgard_block *block,
+                  const midgard_instruction *start, int src)
 {
    mir_compute_liveness(ctx);
 
diff --git a/src/panfrost/midgard/midgard_print.c b/src/panfrost/midgard/midgard_print.c
index 6fe3746ab34..963c0ef2091 100644
--- a/src/panfrost/midgard/midgard_print.c
+++ b/src/panfrost/midgard/midgard_print.c
@@ -79,7 +79,7 @@ mir_print_mask(unsigned mask)
  * don't matter.
  */
 static void
-mir_print_swizzle(unsigned mask, unsigned *swizzle)
+mir_print_swizzle(unsigned mask, const unsigned *swizzle)
 {
    printf(".");
 
@@ -115,7 +115,7 @@ mir_get_unit(unsigned unit)
 }
 
 static void
-mir_print_embedded_constant(midgard_instruction *ins, unsigned src_idx)
+mir_print_embedded_constant(const midgard_instruction *ins, unsigned src_idx)
 {
    assert(src_idx <= 1);
 
@@ -123,7 +123,7 @@ mir_print_embedded_constant(midgard_instruction *ins, unsigned src_idx)
    unsigned sz = nir_alu_type_get_type_size(ins->src_types[src_idx]);
    bool half = (sz == (base_size >> 1));
    unsigned mod = mir_pack_mod(ins, src_idx, false);
-   unsigned *swizzle = ins->swizzle[src_idx];
+   const unsigned *swizzle = ins->swizzle[src_idx];
    midgard_reg_mode reg_mode = reg_mode_for_bitsize(max_bitsize_for_alu(ins));
    unsigned comp_mask = effective_writemask(ins->op, ins->mask);
    unsigned num_comp = util_bitcount(comp_mask);
@@ -153,7 +153,7 @@ mir_print_embedded_constant(midgard_instruction *ins, unsigned src_idx)
 }
 
 static void
-mir_print_src(midgard_instruction *ins, unsigned c)
+mir_print_src(const midgard_instruction *ins, unsigned c)
 {
    mir_print_index(ins->src[c]);
 
@@ -164,7 +164,7 @@ mir_print_src(midgard_instruction *ins, unsigned c)
 }
 
 void
-mir_print_instruction(midgard_instruction *ins)
+mir_print_instruction(const midgard_instruction *ins)
 {
    printf("\t");
 
@@ -326,7 +326,7 @@ mir_print_instruction(midgard_instruction *ins)
 /* Dumps MIR for a block or entire shader respective */
 
 void
-mir_print_block(midgard_block *block)
+mir_print_block(const midgard_block *block)
 {
    printf("block%u: {\n", block->base.name);
 
@@ -360,9 +360,9 @@ mir_print_block(midgard_block *block)
 }
 
 void
-mir_print_shader(compiler_context *ctx)
+mir_print_shader(const compiler_context *ctx)
 {
    mir_foreach_block(ctx, block) {
-      mir_print_block((midgard_block *)block);
+      mir_print_block((const midgard_block *)block);
    }
 }
diff --git a/src/panfrost/midgard/midgard_ra.c b/src/panfrost/midgard/midgard_ra.c
index 599cba53cb2..bf7b3ce304e 100644
--- a/src/panfrost/midgard/midgard_ra.c
+++ b/src/panfrost/midgard/midgard_ra.c
@@ -445,7 +445,7 @@ mir_compute_interference(compiler_context *ctx, struct lcra_state *l)
 }
 
 static bool
-mir_is_64(midgard_instruction *ins)
+mir_is_64(const midgard_instruction *ins)
 {
    if (nir_alu_type_get_type_size(ins->dest_type) == 64)
       return true;
@@ -463,7 +463,7 @@ mir_is_64(midgard_instruction *ins)
  * allocation. TODO: Optimize if barriers and local memory are unused.
  */
 static bool
-needs_contiguous_workgroup(compiler_context *ctx)
+needs_contiguous_workgroup(const compiler_context *ctx)
 {
    return gl_shader_stage_uses_workgroup(ctx->stage);
 }
@@ -475,7 +475,7 @@ needs_contiguous_workgroup(compiler_context *ctx)
  * workgroups.
  */
 static unsigned
-max_threads_per_workgroup(compiler_context *ctx)
+max_threads_per_workgroup(const compiler_context *ctx)
 {
    if (ctx->nir->info.workgroup_size_variable) {
       return 128;
@@ -502,7 +502,7 @@ max_threads_per_workgroup(compiler_context *ctx)
  *    work properly).
  */
 static unsigned
-max_work_registers(compiler_context *ctx)
+max_work_registers(const compiler_context *ctx)
 {
    if (ctx->inputs->is_blend)
       return 8;
diff --git a/src/panfrost/midgard/mir.c b/src/panfrost/midgard/mir.c
index aebf037b0dd..6d06ea9490f 100644
--- a/src/panfrost/midgard/mir.c
+++ b/src/panfrost/midgard/mir.c
@@ -128,7 +128,8 @@ mir_single_use(compiler_context *ctx, unsigned value)
 }
 
 bool
-mir_nontrivial_mod(midgard_instruction *ins, unsigned i, bool check_swizzle)
+mir_nontrivial_mod(const midgard_instruction *ins, unsigned i,
+                   bool check_swizzle)
 {
    bool is_int = midgard_is_integer_op(ins->op);
 
@@ -158,7 +159,7 @@ mir_nontrivial_mod(midgard_instruction *ins, unsigned i, bool check_swizzle)
 }
 
 bool
-mir_nontrivial_outmod(midgard_instruction *ins)
+mir_nontrivial_outmod(const midgard_instruction *ins)
 {
    bool is_int = midgard_is_integer_op(ins->op);
    unsigned mod = ins->outmod;
@@ -232,7 +233,7 @@ mir_round_bytemask_up(uint16_t mask, unsigned bits)
 /* Grabs the per-byte mask of an instruction (as opposed to per-component) */
 
 uint16_t
-mir_bytemask(midgard_instruction *ins)
+mir_bytemask(const midgard_instruction *ins)
 {
    unsigned type_size = nir_alu_type_get_type_size(ins->dest_type);
    return pan_to_bytemask(type_size, ins->mask);
@@ -251,7 +252,7 @@ mir_set_bytemask(midgard_instruction *ins, uint16_t bytemask)
  * for a lower override and negative for no override.
  */
 signed
-mir_upper_override(midgard_instruction *ins, unsigned inst_size)
+mir_upper_override(const midgard_instruction *ins, unsigned inst_size)
 {
    unsigned type_size = nir_alu_type_get_type_size(ins->dest_type);
 
@@ -282,7 +283,7 @@ mir_upper_override(midgard_instruction *ins, unsigned inst_size)
  */
 
 static uint16_t
-mir_bytemask_of_read_components_single(unsigned *swizzle, unsigned inmask,
+mir_bytemask_of_read_components_single(const unsigned *swizzle, unsigned inmask,
                                        unsigned bits)
 {
    unsigned cmask = 0;
@@ -297,7 +298,8 @@ mir_bytemask_of_read_components_single(unsigned *swizzle, unsigned inmask,
 }
 
 uint16_t
-mir_bytemask_of_read_components_index(midgard_instruction *ins, unsigned i)
+mir_bytemask_of_read_components_index(const midgard_instruction *ins,
+                                      unsigned i)
 {
    /* Conditional branches read one 32-bit component = 4 bytes (TODO: multi
     * branch??) */
@@ -326,7 +328,7 @@ mir_bytemask_of_read_components_index(midgard_instruction *ins, unsigned i)
 }
 
 uint16_t
-mir_bytemask_of_read_components(midgard_instruction *ins, unsigned node)
+mir_bytemask_of_read_components(const midgard_instruction *ins, unsigned node)
 {
    uint16_t mask = 0;
 
@@ -374,7 +376,7 @@ mir_bundle_for_op(compiler_context *ctx, midgard_instruction ins)
 }
 
 static unsigned
-mir_bundle_idx_for_ins(midgard_instruction *tag, midgard_block *block)
+mir_bundle_idx_for_ins(const midgard_instruction *tag, midgard_block *block)
 {
    midgard_bundle *bundles = (midgard_bundle *)block->bundles.data;
 
@@ -394,7 +396,7 @@ mir_bundle_idx_for_ins(midgard_instruction *tag, midgard_block *block)
 midgard_instruction *
 mir_insert_instruction_before_scheduled(compiler_context *ctx,
                                         midgard_block *block,
-                                        midgard_instruction *tag,
+                                        const midgard_instruction *tag,
                                         midgard_instruction ins)
 {
    unsigned before = mir_bundle_idx_for_ins(tag, block);
@@ -419,7 +421,7 @@ mir_insert_instruction_before_scheduled(compiler_context *ctx,
 midgard_instruction *
 mir_insert_instruction_after_scheduled(compiler_context *ctx,
                                        midgard_block *block,
-                                       midgard_instruction *tag,
+                                       const midgard_instruction *tag,
                                        midgard_instruction ins)
 {
    /* We need to grow the bundles array to add our new bundle */