diff --git a/src/amd/vulkan/radv_acceleration_structure.c b/src/amd/vulkan/radv_acceleration_structure.c
index c7b6343141d..c236c4ea6d5 100644
--- a/src/amd/vulkan/radv_acceleration_structure.c
+++ b/src/amd/vulkan/radv_acceleration_structure.c
@@ -752,14 +752,13 @@ get_indices(nir_builder *b, nir_ssa_def *addr, nir_ssa_def *type, nir_ssa_def *i
       nir_variable_create(b->shader, nir_var_shader_temp, uvec3_type, "indices");
 
    nir_push_if(b, nir_ult(b, type, nir_imm_int(b, 2)));
-   nir_push_if(b, nir_ieq(b, type, nir_imm_int(b, VK_INDEX_TYPE_UINT16)));
+   nir_push_if(b, nir_ieq_imm(b, type, VK_INDEX_TYPE_UINT16));
    {
       nir_ssa_def *index_id = nir_umul24(b, id, nir_imm_int(b, 6));
       nir_ssa_def *indices[3];
       for (unsigned i = 0; i < 3; ++i) {
          indices[i] = nir_build_load_global(
-            b, 1, 16,
-            nir_iadd(b, addr, nir_u2u64(b, nir_iadd(b, index_id, nir_imm_int(b, 2 * i)))));
+            b, 1, 16, nir_iadd(b, addr, nir_u2u64(b, nir_iadd_imm(b, index_id, 2 * i))));
       }
       nir_store_var(b, result, nir_u2u32(b, nir_vec(b, indices, 3)), 7);
    }
@@ -776,11 +775,11 @@ get_indices(nir_builder *b, nir_ssa_def *addr, nir_ssa_def *type, nir_ssa_def *i
       nir_ssa_def *index_id = nir_umul24(b, id, nir_imm_int(b, 3));
       nir_ssa_def *indices[] = {
          index_id,
-         nir_iadd(b, index_id, nir_imm_int(b, 1)),
-         nir_iadd(b, index_id, nir_imm_int(b, 2)),
+         nir_iadd_imm(b, index_id, 1),
+         nir_iadd_imm(b, index_id, 2),
       };
 
-      nir_push_if(b, nir_ieq(b, type, nir_imm_int(b, VK_INDEX_TYPE_NONE_KHR)));
+      nir_push_if(b, nir_ieq_imm(b, type, VK_INDEX_TYPE_NONE_KHR));
       {
          nir_store_var(b, result, nir_vec(b, indices, 3), 7);
       }
@@ -827,7 +826,7 @@ get_vertices(nir_builder *b, nir_ssa_def *addresses, nir_ssa_def *format, nir_ss
 
    for (unsigned f = 0; f < ARRAY_SIZE(formats); ++f) {
       if (f + 1 < ARRAY_SIZE(formats))
-         nir_push_if(b, nir_ieq(b, format, nir_imm_int(b, formats[f])));
+         nir_push_if(b, nir_ieq_imm(b, format, formats[f]));
 
       for (unsigned i = 0; i < 3; ++i) {
          switch (formats[f]) {
@@ -863,8 +862,8 @@ get_vertices(nir_builder *b, nir_ssa_def *addresses, nir_ssa_def *format, nir_ss
                   values[j] = nir_ubfe(b, val, nir_imm_int(b, j * 10), nir_imm_int(b, 10));
             } else {
                for (unsigned j = 0; j < components; ++j)
-                  values[j] = nir_build_load_global(
-                     b, 1, comp_bits, nir_iadd(b, addr, nir_imm_int64(b, j * comp_bytes)));
+                  values[j] =
+                     nir_build_load_global(b, 1, comp_bits, nir_iadd_imm(b, addr, j * comp_bytes));
 
                for (unsigned j = components; j < 3; ++j)
                   values[j] = nir_imm_intN_t(b, 0, comp_bits);
@@ -1007,7 +1006,7 @@ build_leaf_shader(struct radv_device *dev)
       &b, scratch_addr,
       nir_u2u64(&b, nir_iadd(&b, scratch_offset, nir_umul24(&b, global_id, nir_imm_int(&b, 4)))));
 
-   nir_push_if(&b, nir_ieq(&b, geom_type, nir_imm_int(&b, VK_GEOMETRY_TYPE_TRIANGLES_KHR)));
+   nir_push_if(&b, nir_ieq_imm(&b, geom_type, VK_GEOMETRY_TYPE_TRIANGLES_KHR));
    { /* Triangles */
       nir_ssa_def *vertex_addr = nir_pack_64_2x32(&b, nir_channels(&b, pconst2, 3));
       nir_ssa_def *index_addr = nir_pack_64_2x32(&b, nir_channels(&b, pconst2, 12));
@@ -1040,18 +1039,13 @@ build_leaf_shader(struct radv_device *dev)
       nir_store_var(&b, transform[1], nir_imm_vec4(&b, 0.0, 1.0, 0.0, 0.0), 0xf);
       nir_store_var(&b, transform[2], nir_imm_vec4(&b, 0.0, 0.0, 1.0, 0.0), 0xf);
 
-      nir_push_if(&b, nir_ine(&b, transform_addr, nir_imm_int64(&b, 0)));
-      nir_store_var(
-         &b, transform[0],
-         nir_build_load_global(&b, 4, 32, nir_iadd(&b, transform_addr, nir_imm_int64(&b, 0))), 0xf);
-      nir_store_var(
-         &b, transform[1],
-         nir_build_load_global(&b, 4, 32, nir_iadd(&b, transform_addr, nir_imm_int64(&b, 16))),
-         0xf);
-      nir_store_var(
-         &b, transform[2],
-         nir_build_load_global(&b, 4, 32, nir_iadd(&b, transform_addr, nir_imm_int64(&b, 32))),
-         0xf);
+      nir_push_if(&b, nir_ine_imm(&b, transform_addr, 0));
+      nir_store_var(&b, transform[0],
+                    nir_build_load_global(&b, 4, 32, nir_iadd_imm(&b, transform_addr, 0)), 0xf);
+      nir_store_var(&b, transform[1],
+                    nir_build_load_global(&b, 4, 32, nir_iadd_imm(&b, transform_addr, 16)), 0xf);
+      nir_store_var(&b, transform[2],
+                    nir_build_load_global(&b, 4, 32, nir_iadd_imm(&b, transform_addr, 32)), 0xf);
       nir_pop_if(&b, NULL);
 
       for (unsigned i = 0; i < 3; ++i)
@@ -1067,15 +1061,14 @@ build_leaf_shader(struct radv_device *dev)
 
       for (unsigned i = 0; i < 4; ++i) {
          nir_build_store_global(&b, nir_vec(&b, node_data + i * 4, 4),
-                                nir_iadd(&b, triangle_node_dst_addr, nir_imm_int64(&b, i * 16)),
-                                .align_mul = 16);
+                                nir_iadd_imm(&b, triangle_node_dst_addr, i * 16), .align_mul = 16);
       }
 
-      nir_ssa_def *node_id = nir_ushr(&b, node_offset, nir_imm_int(&b, 3));
+      nir_ssa_def *node_id = nir_ushr_imm(&b, node_offset, 3);
       nir_build_store_global(&b, node_id, scratch_addr);
    }
    nir_push_else(&b, NULL);
-   nir_push_if(&b, nir_ieq(&b, geom_type, nir_imm_int(&b, VK_GEOMETRY_TYPE_AABBS_KHR)));
+   nir_push_if(&b, nir_ieq_imm(&b, geom_type, VK_GEOMETRY_TYPE_AABBS_KHR));
    { /* AABBs */
       nir_ssa_def *aabb_addr = nir_pack_64_2x32(&b, nir_channels(&b, pconst2, 3));
       nir_ssa_def *aabb_stride = nir_channel(&b, pconst2, 2);
@@ -1083,16 +1076,13 @@ build_leaf_shader(struct radv_device *dev)
       nir_ssa_def *node_offset =
          nir_iadd(&b, node_dst_offset, nir_umul24(&b, global_id, nir_imm_int(&b, 64)));
       nir_ssa_def *aabb_node_dst_addr = nir_iadd(&b, node_dst_addr, nir_u2u64(&b, node_offset));
-      nir_ssa_def *node_id =
-         nir_iadd(&b, nir_ushr(&b, node_offset, nir_imm_int(&b, 3)), nir_imm_int(&b, 7));
+      nir_ssa_def *node_id = nir_iadd_imm(&b, nir_ushr_imm(&b, node_offset, 3), 7);
       nir_build_store_global(&b, node_id, scratch_addr);
 
       aabb_addr = nir_iadd(&b, aabb_addr, nir_u2u64(&b, nir_imul(&b, aabb_stride, global_id)));
 
-      nir_ssa_def *min_bound =
-         nir_build_load_global(&b, 3, 32, nir_iadd(&b, aabb_addr, nir_imm_int64(&b, 0)));
-      nir_ssa_def *max_bound =
-         nir_build_load_global(&b, 3, 32, nir_iadd(&b, aabb_addr, nir_imm_int64(&b, 12)));
+      nir_ssa_def *min_bound = nir_build_load_global(&b, 3, 32, nir_iadd_imm(&b, aabb_addr, 0));
+      nir_ssa_def *max_bound = nir_build_load_global(&b, 3, 32, nir_iadd_imm(&b, aabb_addr, 12));
 
       nir_ssa_def *values[] = {nir_channel(&b, min_bound, 0),
                                nir_channel(&b, min_bound, 1),
@@ -1104,21 +1094,19 @@ build_leaf_shader(struct radv_device *dev)
                                geometry_id};
 
       nir_build_store_global(&b, nir_vec(&b, values + 0, 4),
-                             nir_iadd(&b, aabb_node_dst_addr, nir_imm_int64(&b, 0)),
-                             .align_mul = 16);
+                             nir_iadd_imm(&b, aabb_node_dst_addr, 0), .align_mul = 16);
       nir_build_store_global(&b, nir_vec(&b, values + 4, 4),
-                             nir_iadd(&b, aabb_node_dst_addr, nir_imm_int64(&b, 16)),
-                             .align_mul = 16);
+                             nir_iadd_imm(&b, aabb_node_dst_addr, 16), .align_mul = 16);
    }
    nir_push_else(&b, NULL);
    { /* Instances */
 
       nir_variable *instance_addr_var =
          nir_variable_create(b.shader, nir_var_shader_temp, glsl_uint64_t_type(), "instance_addr");
-      nir_push_if(&b, nir_ine(&b, nir_channel(&b, pconst2, 2), nir_imm_int(&b, 0)));
+      nir_push_if(&b, nir_ine_imm(&b, nir_channel(&b, pconst2, 2), 0));
       {
          nir_ssa_def *ptr = nir_iadd(&b, nir_pack_64_2x32(&b, nir_channels(&b, pconst2, 3)),
-                                     nir_u2u64(&b, nir_imul(&b, global_id, nir_imm_int(&b, 8))));
+                                     nir_u2u64(&b, nir_imul_imm(&b, global_id, 8)));
          nir_ssa_def *addr =
             nir_pack_64_2x32(&b, nir_build_load_global(&b, 2, 32, ptr, .align_mul = 8));
          nir_store_var(&b, instance_addr_var, addr, 1);
@@ -1126,24 +1114,22 @@ build_leaf_shader(struct radv_device *dev)
       nir_push_else(&b, NULL);
       {
          nir_ssa_def *addr = nir_iadd(&b, nir_pack_64_2x32(&b, nir_channels(&b, pconst2, 3)),
-                                      nir_u2u64(&b, nir_imul(&b, global_id, nir_imm_int(&b, 64))));
+                                      nir_u2u64(&b, nir_imul_imm(&b, global_id, 64)));
          nir_store_var(&b, instance_addr_var, addr, 1);
       }
       nir_pop_if(&b, NULL);
       nir_ssa_def *instance_addr = nir_load_var(&b, instance_addr_var);
 
       nir_ssa_def *inst_transform[] = {
-         nir_build_load_global(&b, 4, 32, nir_iadd(&b, instance_addr, nir_imm_int64(&b, 0))),
-         nir_build_load_global(&b, 4, 32, nir_iadd(&b, instance_addr, nir_imm_int64(&b, 16))),
-         nir_build_load_global(&b, 4, 32, nir_iadd(&b, instance_addr, nir_imm_int64(&b, 32)))};
-      nir_ssa_def *inst3 =
-         nir_build_load_global(&b, 4, 32, nir_iadd(&b, instance_addr, nir_imm_int64(&b, 48)));
+         nir_build_load_global(&b, 4, 32, nir_iadd_imm(&b, instance_addr, 0)),
+         nir_build_load_global(&b, 4, 32, nir_iadd_imm(&b, instance_addr, 16)),
+         nir_build_load_global(&b, 4, 32, nir_iadd_imm(&b, instance_addr, 32))};
+      nir_ssa_def *inst3 = nir_build_load_global(&b, 4, 32, nir_iadd_imm(&b, instance_addr, 48));
 
       nir_ssa_def *node_offset =
          nir_iadd(&b, node_dst_offset, nir_umul24(&b, global_id, nir_imm_int(&b, 128)));
       node_dst_addr = nir_iadd(&b, node_dst_addr, nir_u2u64(&b, node_offset));
-      nir_ssa_def *node_id =
-         nir_iadd(&b, nir_ushr(&b, node_offset, nir_imm_int(&b, 3)), nir_imm_int(&b, 6));
+      nir_ssa_def *node_id = nir_iadd_imm(&b, nir_ushr_imm(&b, node_offset, 3), 6);
       nir_build_store_global(&b, node_id, scratch_addr);
 
       nir_variable *bounds[2] = {
@@ -1155,13 +1141,11 @@ build_leaf_shader(struct radv_device *dev)
       nir_store_var(&b, bounds[1], nir_channels(&b, nir_imm_vec4(&b, NAN, NAN, NAN, NAN), 7), 7);
 
       nir_ssa_def *header_addr = nir_pack_64_2x32(&b, nir_channels(&b, inst3, 12));
-      nir_push_if(&b, nir_ine(&b, header_addr, nir_imm_int64(&b, 0)));
+      nir_push_if(&b, nir_ine_imm(&b, header_addr, 0));
       nir_ssa_def *header_root_offset =
-         nir_build_load_global(&b, 1, 32, nir_iadd(&b, header_addr, nir_imm_int64(&b, 0)));
-      nir_ssa_def *header_min =
-         nir_build_load_global(&b, 3, 32, nir_iadd(&b, header_addr, nir_imm_int64(&b, 8)));
-      nir_ssa_def *header_max =
-         nir_build_load_global(&b, 3, 32, nir_iadd(&b, header_addr, nir_imm_int64(&b, 20)));
+         nir_build_load_global(&b, 1, 32, nir_iadd_imm(&b, header_addr, 0));
+      nir_ssa_def *header_min = nir_build_load_global(&b, 3, 32, nir_iadd_imm(&b, header_addr, 8));
+      nir_ssa_def *header_max = nir_build_load_global(&b, 3, 32, nir_iadd_imm(&b, header_addr, 20));
 
       nir_ssa_def *bound_defs[2][3];
       for (unsigned i = 0; i < 3; ++i) {
@@ -1187,7 +1171,7 @@ build_leaf_shader(struct radv_device *dev)
             vals[j] = nir_channel(&b, inst_transform[j], i);
 
          nir_build_store_global(&b, nir_vec(&b, vals, 3),
-                                nir_iadd(&b, node_dst_addr, nir_imm_int64(&b, 92 + 12 * i)));
+                                nir_iadd_imm(&b, node_dst_addr, 92 + 12 * i));
       }
 
       nir_ssa_def *m_in[3][3], *m_out[3][3], *m_vec[3][4];
@@ -1203,21 +1187,18 @@ build_leaf_shader(struct radv_device *dev)
 
       for (unsigned i = 0; i < 3; ++i) {
          nir_build_store_global(&b, nir_vec(&b, m_vec[i], 4),
-                                nir_iadd(&b, node_dst_addr, nir_imm_int64(&b, 16 + 16 * i)));
+                                nir_iadd_imm(&b, node_dst_addr, 16 + 16 * i));
       }
 
       nir_ssa_def *out0[4] = {
          nir_ior(&b, nir_channel(&b, nir_unpack_64_2x32(&b, header_addr), 0), header_root_offset),
          nir_channel(&b, nir_unpack_64_2x32(&b, header_addr), 1), nir_channel(&b, inst3, 0),
          nir_channel(&b, inst3, 1)};
-      nir_build_store_global(&b, nir_vec(&b, out0, 4),
-                             nir_iadd(&b, node_dst_addr, nir_imm_int64(&b, 0)));
-      nir_build_store_global(&b, global_id, nir_iadd(&b, node_dst_addr, nir_imm_int64(&b, 88)));
+      nir_build_store_global(&b, nir_vec(&b, out0, 4), nir_iadd_imm(&b, node_dst_addr, 0));
+      nir_build_store_global(&b, global_id, nir_iadd_imm(&b, node_dst_addr, 88));
       nir_pop_if(&b, NULL);
-      nir_build_store_global(&b, nir_load_var(&b, bounds[0]),
-                             nir_iadd(&b, node_dst_addr, nir_imm_int64(&b, 64)));
-      nir_build_store_global(&b, nir_load_var(&b, bounds[1]),
-                             nir_iadd(&b, node_dst_addr, nir_imm_int64(&b, 76)));
+      nir_build_store_global(&b, nir_load_var(&b, bounds[0]), nir_iadd_imm(&b, node_dst_addr, 64));
+      nir_build_store_global(&b, nir_load_var(&b, bounds[1]), nir_iadd_imm(&b, node_dst_addr, 76));
    }
    nir_pop_if(&b, NULL);
    nir_pop_if(&b, NULL);
@@ -1229,17 +1210,15 @@ static void
 determine_bounds(nir_builder *b, nir_ssa_def *node_addr, nir_ssa_def *node_id,
                  nir_variable *bounds_vars[2])
 {
-   nir_ssa_def *node_type = nir_iand(b, node_id, nir_imm_int(b, 7));
-   node_addr = nir_iadd(
-      b, node_addr,
-      nir_u2u64(b, nir_ishl(b, nir_iand(b, node_id, nir_imm_int(b, ~7u)), nir_imm_int(b, 3))));
+   nir_ssa_def *node_type = nir_iand_imm(b, node_id, 7);
+   node_addr =
+      nir_iadd(b, node_addr, nir_u2u64(b, nir_ishl_imm(b, nir_iand_imm(b, node_id, ~7u), 3)));
 
-   nir_push_if(b, nir_ieq(b, node_type, nir_imm_int(b, 0)));
+   nir_push_if(b, nir_ieq_imm(b, node_type, 0));
    {
       nir_ssa_def *positions[3];
       for (unsigned i = 0; i < 3; ++i)
-         positions[i] =
-            nir_build_load_global(b, 3, 32, nir_iadd(b, node_addr, nir_imm_int64(b, i * 12)));
+         positions[i] = nir_build_load_global(b, 3, 32, nir_iadd_imm(b, node_addr, i * 12));
       nir_ssa_def *bounds[] = {positions[0], positions[0]};
       for (unsigned i = 1; i < 3; ++i) {
          bounds[0] = nir_fmin(b, bounds[0], positions[i]);
@@ -1249,13 +1228,13 @@ determine_bounds(nir_builder *b, nir_ssa_def *node_addr, nir_ssa_def *node_id,
       nir_store_var(b, bounds_vars[1], bounds[1], 7);
    }
    nir_push_else(b, NULL);
-   nir_push_if(b, nir_ieq(b, node_type, nir_imm_int(b, 5)));
+   nir_push_if(b, nir_ieq_imm(b, node_type, 5));
    {
       nir_ssa_def *input_bounds[4][2];
       for (unsigned i = 0; i < 4; ++i)
          for (unsigned j = 0; j < 2; ++j)
-            input_bounds[i][j] = nir_build_load_global(
-               b, 3, 32, nir_iadd(b, node_addr, nir_imm_int64(b, 16 + i * 24 + j * 12)));
+            input_bounds[i][j] =
+               nir_build_load_global(b, 3, 32, nir_iadd_imm(b, node_addr, 16 + i * 24 + j * 12));
       nir_ssa_def *bounds[] = {input_bounds[0][0], input_bounds[0][1]};
       for (unsigned i = 1; i < 4; ++i) {
          bounds[0] = nir_fmin(b, bounds[0], input_bounds[i][0]);
@@ -1266,12 +1245,11 @@ determine_bounds(nir_builder *b, nir_ssa_def *node_addr, nir_ssa_def *node_id,
       nir_store_var(b, bounds_vars[1], bounds[1], 7);
    }
    nir_push_else(b, NULL);
-   nir_push_if(b, nir_ieq(b, node_type, nir_imm_int(b, 6)));
+   nir_push_if(b, nir_ieq_imm(b, node_type, 6));
    { /* Instances */
       nir_ssa_def *bounds[2];
       for (unsigned i = 0; i < 2; ++i)
-         bounds[i] =
-            nir_build_load_global(b, 3, 32, nir_iadd(b, node_addr, nir_imm_int64(b, 64 + i * 12)));
+         bounds[i] = nir_build_load_global(b, 3, 32, nir_iadd_imm(b, node_addr, 64 + i * 12));
       nir_store_var(b, bounds_vars[0], bounds[0], 7);
       nir_store_var(b, bounds_vars[1], bounds[1], 7);
    }
@@ -1279,8 +1257,7 @@ determine_bounds(nir_builder *b, nir_ssa_def *node_addr, nir_ssa_def *node_id,
    { /* AABBs */
       nir_ssa_def *bounds[2];
       for (unsigned i = 0; i < 2; ++i)
-         bounds[i] =
-            nir_build_load_global(b, 3, 32, nir_iadd(b, node_addr, nir_imm_int64(b, i * 12)));
+         bounds[i] = nir_build_load_global(b, 3, 32, nir_iadd_imm(b, node_addr, i * 12));
       nir_store_var(b, bounds_vars[0], bounds[0], 7);
       nir_store_var(b, bounds_vars[1], bounds[1], 7);
    }
@@ -1316,30 +1293,26 @@ build_internal_shader(struct radv_device *dev)
    nir_ssa_def *node_dst_offset = nir_channel(&b, pconst1, 0);
    nir_ssa_def *dst_scratch_offset = nir_channel(&b, pconst1, 1);
    nir_ssa_def *src_scratch_offset = nir_channel(&b, pconst1, 2);
-   nir_ssa_def *src_node_count =
-      nir_iand(&b, nir_channel(&b, pconst1, 3), nir_imm_int(&b, 0x7FFFFFFFU));
+   nir_ssa_def *src_node_count = nir_iand_imm(&b, nir_channel(&b, pconst1, 3), 0x7FFFFFFFU);
    nir_ssa_def *fill_header =
-      nir_ine(&b, nir_iand(&b, nir_channel(&b, pconst1, 3), nir_imm_int(&b, 0x80000000U)),
-              nir_imm_int(&b, 0));
+      nir_ine_imm(&b, nir_iand_imm(&b, nir_channel(&b, pconst1, 3), 0x80000000U), 0);
 
    nir_ssa_def *global_id =
       nir_iadd(&b,
                nir_umul24(&b, nir_channels(&b, nir_load_workgroup_id(&b, 32), 1),
                           nir_imm_int(&b, b.shader->info.workgroup_size[0])),
                nir_channels(&b, nir_load_local_invocation_id(&b), 1));
-   nir_ssa_def *src_idx = nir_imul(&b, global_id, nir_imm_int(&b, 4));
+   nir_ssa_def *src_idx = nir_imul_imm(&b, global_id, 4);
    nir_ssa_def *src_count = nir_umin(&b, nir_imm_int(&b, 4), nir_isub(&b, src_node_count, src_idx));
 
-   nir_ssa_def *node_offset =
-      nir_iadd(&b, node_dst_offset, nir_ishl(&b, global_id, nir_imm_int(&b, 7)));
+   nir_ssa_def *node_offset = nir_iadd(&b, node_dst_offset, nir_ishl_imm(&b, global_id, 7));
    nir_ssa_def *node_dst_addr = nir_iadd(&b, node_addr, nir_u2u64(&b, node_offset));
    nir_ssa_def *src_nodes = nir_build_load_global(
       &b, 4, 32,
       nir_iadd(&b, scratch_addr,
-               nir_u2u64(&b, nir_iadd(&b, src_scratch_offset,
-                                      nir_ishl(&b, global_id, nir_imm_int(&b, 4))))));
+               nir_u2u64(&b, nir_iadd(&b, src_scratch_offset, nir_ishl_imm(&b, global_id, 4)))));
 
-   nir_build_store_global(&b, src_nodes, nir_iadd(&b, node_dst_addr, nir_imm_int64(&b, 0)));
+   nir_build_store_global(&b, src_nodes, nir_iadd_imm(&b, node_dst_addr, 0));
 
    nir_ssa_def *total_bounds[2] = {
       nir_channels(&b, nir_imm_vec4(&b, NAN, NAN, NAN, NAN), 7),
@@ -1358,24 +1331,23 @@ build_internal_shader(struct radv_device *dev)
       determine_bounds(&b, node_addr, nir_channel(&b, src_nodes, i), bounds);
       nir_pop_if(&b, NULL);
       nir_build_store_global(&b, nir_load_var(&b, bounds[0]),
-                             nir_iadd(&b, node_dst_addr, nir_imm_int64(&b, 16 + 24 * i)));
+                             nir_iadd_imm(&b, node_dst_addr, 16 + 24 * i));
       nir_build_store_global(&b, nir_load_var(&b, bounds[1]),
-                             nir_iadd(&b, node_dst_addr, nir_imm_int64(&b, 28 + 24 * i)));
+                             nir_iadd_imm(&b, node_dst_addr, 28 + 24 * i));
       total_bounds[0] = nir_fmin(&b, total_bounds[0], nir_load_var(&b, bounds[0]));
       total_bounds[1] = nir_fmax(&b, total_bounds[1], nir_load_var(&b, bounds[1]));
    }
 
-   nir_ssa_def *node_id =
-      nir_iadd(&b, nir_ushr(&b, node_offset, nir_imm_int(&b, 3)), nir_imm_int(&b, 5));
-   nir_ssa_def *dst_scratch_addr = nir_iadd(
-      &b, scratch_addr,
-      nir_u2u64(&b, nir_iadd(&b, dst_scratch_offset, nir_ishl(&b, global_id, nir_imm_int(&b, 2)))));
+   nir_ssa_def *node_id = nir_iadd_imm(&b, nir_ushr_imm(&b, node_offset, 3), 5);
+   nir_ssa_def *dst_scratch_addr =
+      nir_iadd(&b, scratch_addr,
+               nir_u2u64(&b, nir_iadd(&b, dst_scratch_offset, nir_ishl_imm(&b, global_id, 2))));
    nir_build_store_global(&b, node_id, dst_scratch_addr);
 
    nir_push_if(&b, fill_header);
    nir_build_store_global(&b, node_id, node_addr);
-   nir_build_store_global(&b, total_bounds[0], nir_iadd(&b, node_addr, nir_imm_int64(&b, 8)));
-   nir_build_store_global(&b, total_bounds[1], nir_iadd(&b, node_addr, nir_imm_int64(&b, 20)));
+   nir_build_store_global(&b, total_bounds[0], nir_iadd_imm(&b, node_addr, 8));
+   nir_build_store_global(&b, total_bounds[1], nir_iadd_imm(&b, node_addr, 20));
    nir_pop_if(&b, NULL);
    return b.shader;
 }
@@ -1409,11 +1381,11 @@ build_copy_shader(struct radv_device *dev)
 
    nir_variable *offset_var =
       nir_variable_create(b.shader, nir_var_shader_temp, glsl_uint_type(), "offset");
-   nir_ssa_def *offset = nir_imul(&b, global_id, nir_imm_int(&b, 16));
+   nir_ssa_def *offset = nir_imul_imm(&b, global_id, 16);
    nir_store_var(&b, offset_var, offset, 1);
 
-   nir_ssa_def *increment = nir_imul(&b, nir_channel(&b, nir_load_num_workgroups(&b, 32), 0),
-                                     nir_imm_int(&b, b.shader->info.workgroup_size[0] * 16));
+   nir_ssa_def *increment = nir_imul_imm(&b, nir_channel(&b, nir_load_num_workgroups(&b, 32), 0),
+                                         b.shader->info.workgroup_size[0] * 16);
 
    nir_ssa_def *pconst0 =
       nir_load_push_constant(&b, 4, 32, nir_imm_int(&b, 0), .base = 0, .range = 16);
@@ -1436,98 +1408,87 @@ build_copy_shader(struct radv_device *dev)
    nir_variable *value_var =
       nir_variable_create(b.shader, nir_var_shader_temp, glsl_vec4_type(), "value");
 
-   nir_push_if(&b, nir_ieq(&b, mode, nir_imm_int(&b, COPY_MODE_SERIALIZE)));
+   nir_push_if(&b, nir_ieq_imm(&b, mode, COPY_MODE_SERIALIZE));
    {
       nir_ssa_def *instance_count = nir_build_load_global(
          &b, 1, 32,
-         nir_iadd(&b, src_base_addr,
-                  nir_imm_int64(&b, offsetof(struct radv_accel_struct_header, instance_count))));
+         nir_iadd_imm(&b, src_base_addr,
+                      offsetof(struct radv_accel_struct_header, instance_count)));
       nir_ssa_def *compacted_size = nir_build_load_global(
          &b, 1, 64,
-         nir_iadd(&b, src_base_addr,
-                  nir_imm_int64(&b, offsetof(struct radv_accel_struct_header, compacted_size))));
+         nir_iadd_imm(&b, src_base_addr,
+                      offsetof(struct radv_accel_struct_header, compacted_size)));
       nir_ssa_def *serialization_size = nir_build_load_global(
          &b, 1, 64,
-         nir_iadd(
-            &b, src_base_addr,
-            nir_imm_int64(&b, offsetof(struct radv_accel_struct_header, serialization_size))));
+         nir_iadd_imm(&b, src_base_addr,
+                      offsetof(struct radv_accel_struct_header, serialization_size)));
 
       nir_store_var(&b, compacted_size_var, compacted_size, 1);
-      nir_store_var(
-         &b, instance_offset_var,
-         nir_build_load_global(&b, 1, 32,
-                               nir_iadd(&b, src_base_addr,
-                                        nir_imm_int64(&b, offsetof(struct radv_accel_struct_header,
-                                                                   instance_offset)))),
-         1);
+      nir_store_var(&b, instance_offset_var,
+                    nir_build_load_global(
+                       &b, 1, 32,
+                       nir_iadd_imm(&b, src_base_addr,
+                                    offsetof(struct radv_accel_struct_header, instance_offset))),
+                    1);
       nir_store_var(&b, instance_count_var, instance_count, 1);
 
-      nir_ssa_def *dst_offset =
-         nir_iadd(&b, nir_imm_int(&b, sizeof(struct radv_accel_struct_serialization_header)),
-                  nir_imul(&b, instance_count, nir_imm_int(&b, sizeof(uint64_t))));
+      nir_ssa_def *dst_offset = nir_iadd_imm(&b, nir_imul_imm(&b, instance_count, sizeof(uint64_t)),
+                                             sizeof(struct radv_accel_struct_serialization_header));
       nir_store_var(&b, src_offset_var, nir_imm_int(&b, 0), 1);
       nir_store_var(&b, dst_offset_var, dst_offset, 1);
 
-      nir_push_if(&b, nir_ieq(&b, global_id, nir_imm_int(&b, 0)));
+      nir_push_if(&b, nir_ieq_imm(&b, global_id, 0));
       {
-         nir_build_store_global(
-            &b, serialization_size,
-            nir_iadd(&b, dst_base_addr,
-                     nir_imm_int64(&b, offsetof(struct radv_accel_struct_serialization_header,
-                                                serialization_size))));
+         nir_build_store_global(&b, serialization_size,
+                                nir_iadd_imm(&b, dst_base_addr,
+                                             offsetof(struct radv_accel_struct_serialization_header,
+                                                      serialization_size)));
          nir_build_store_global(
             &b, compacted_size,
-            nir_iadd(&b, dst_base_addr,
-                     nir_imm_int64(&b, offsetof(struct radv_accel_struct_serialization_header,
-                                                compacted_size))));
+            nir_iadd_imm(&b, dst_base_addr,
+                         offsetof(struct radv_accel_struct_serialization_header, compacted_size)));
          nir_build_store_global(
             &b, nir_u2u64(&b, instance_count),
-            nir_iadd(&b, dst_base_addr,
-                     nir_imm_int64(&b, offsetof(struct radv_accel_struct_serialization_header,
-                                                instance_count))));
+            nir_iadd_imm(&b, dst_base_addr,
+                         offsetof(struct radv_accel_struct_serialization_header, instance_count)));
       }
       nir_pop_if(&b, NULL);
    }
    nir_push_else(&b, NULL);
-   nir_push_if(&b, nir_ieq(&b, mode, nir_imm_int(&b, COPY_MODE_DESERIALIZE)));
+   nir_push_if(&b, nir_ieq_imm(&b, mode, COPY_MODE_DESERIALIZE));
    {
       nir_ssa_def *instance_count = nir_build_load_global(
          &b, 1, 32,
-         nir_iadd(&b, src_base_addr,
-                  nir_imm_int64(
-                     &b, offsetof(struct radv_accel_struct_serialization_header, instance_count))));
-      nir_ssa_def *src_offset =
-         nir_iadd(&b, nir_imm_int(&b, sizeof(struct radv_accel_struct_serialization_header)),
-                  nir_imul(&b, instance_count, nir_imm_int(&b, sizeof(uint64_t))));
+         nir_iadd_imm(&b, src_base_addr,
+                      offsetof(struct radv_accel_struct_serialization_header, instance_count)));
+      nir_ssa_def *src_offset = nir_iadd_imm(&b, nir_imul_imm(&b, instance_count, sizeof(uint64_t)),
+                                             sizeof(struct radv_accel_struct_serialization_header));
 
       nir_ssa_def *header_addr = nir_iadd(&b, src_base_addr, nir_u2u64(&b, src_offset));
-      nir_store_var(
-         &b, compacted_size_var,
-         nir_build_load_global(
-            &b, 1, 64,
-            nir_iadd(&b, header_addr,
-                     nir_imm_int64(&b, offsetof(struct radv_accel_struct_header, compacted_size)))),
-         1);
-      nir_store_var(
-         &b, instance_offset_var,
-         nir_build_load_global(&b, 1, 32,
-                               nir_iadd(&b, header_addr,
-                                        nir_imm_int64(&b, offsetof(struct radv_accel_struct_header,
-                                                                   instance_offset)))),
-         1);
+      nir_store_var(&b, compacted_size_var,
+                    nir_build_load_global(
+                       &b, 1, 64,
+                       nir_iadd_imm(&b, header_addr,
+                                    offsetof(struct radv_accel_struct_header, compacted_size))),
+                    1);
+      nir_store_var(&b, instance_offset_var,
+                    nir_build_load_global(
+                       &b, 1, 32,
+                       nir_iadd_imm(&b, header_addr,
+                                    offsetof(struct radv_accel_struct_header, instance_offset))),
+                    1);
       nir_store_var(&b, instance_count_var, instance_count, 1);
       nir_store_var(&b, src_offset_var, src_offset, 1);
       nir_store_var(&b, dst_offset_var, nir_imm_int(&b, 0), 1);
    }
    nir_push_else(&b, NULL); /* COPY_MODE_COPY */
    {
-      nir_store_var(
-         &b, compacted_size_var,
-         nir_build_load_global(
-            &b, 1, 64,
-            nir_iadd(&b, src_base_addr,
-                     nir_imm_int64(&b, offsetof(struct radv_accel_struct_header, compacted_size)))),
-         1);
+      nir_store_var(&b, compacted_size_var,
+                    nir_build_load_global(
+                       &b, 1, 64,
+                       nir_iadd_imm(&b, src_base_addr,
+                                    offsetof(struct radv_accel_struct_header, compacted_size))),
+                    1);
 
       nir_store_var(&b, src_offset_var, nir_imm_int(&b, 0), 1);
       nir_store_var(&b, dst_offset_var, nir_imm_int(&b, 0), 1);
@@ -1538,12 +1499,10 @@ build_copy_shader(struct radv_device *dev)
    nir_pop_if(&b, NULL);
 
    nir_ssa_def *instance_bound =
-      nir_imul(&b, nir_imm_int(&b, sizeof(struct radv_bvh_instance_node)),
-               nir_load_var(&b, instance_count_var));
+      nir_imul_imm(&b, nir_load_var(&b, instance_count_var), sizeof(struct radv_bvh_instance_node));
    nir_ssa_def *compacted_size = nir_build_load_global(
       &b, 1, 32,
-      nir_iadd(&b, src_base_addr,
-               nir_imm_int64(&b, offsetof(struct radv_accel_struct_header, compacted_size))));
+      nir_iadd_imm(&b, src_base_addr, offsetof(struct radv_accel_struct_header, compacted_size)));
 
    nir_push_loop(&b);
    {
@@ -1562,23 +1521,18 @@ build_copy_shader(struct radv_device *dev)
          nir_ssa_def *in_instance_bound =
             nir_iand(&b, nir_uge(&b, offset, nir_load_var(&b, instance_offset_var)),
                      nir_ult(&b, instance_offset, instance_bound));
-         nir_ssa_def *instance_start =
-            nir_ieq(&b,
-                    nir_iand(&b, instance_offset,
-                             nir_imm_int(&b, sizeof(struct radv_bvh_instance_node) - 1)),
-                    nir_imm_int(&b, 0));
+         nir_ssa_def *instance_start = nir_ieq_imm(
+            &b, nir_iand_imm(&b, instance_offset, sizeof(struct radv_bvh_instance_node) - 1), 0);
 
          nir_push_if(&b, nir_iand(&b, in_instance_bound, instance_start));
          {
-            nir_ssa_def *instance_id = nir_ushr(&b, instance_offset, nir_imm_int(&b, 7));
+            nir_ssa_def *instance_id = nir_ushr_imm(&b, instance_offset, 7);
 
-            nir_push_if(&b, nir_ieq(&b, mode, nir_imm_int(&b, COPY_MODE_SERIALIZE)));
+            nir_push_if(&b, nir_ieq_imm(&b, mode, COPY_MODE_SERIALIZE));
             {
-               nir_ssa_def *instance_addr =
-                  nir_imul(&b, instance_id, nir_imm_int(&b, sizeof(uint64_t)));
-               instance_addr =
-                  nir_iadd(&b, instance_addr,
-                           nir_imm_int(&b, sizeof(struct radv_accel_struct_serialization_header)));
+               nir_ssa_def *instance_addr = nir_imul_imm(&b, instance_id, sizeof(uint64_t));
+               instance_addr = nir_iadd_imm(&b, instance_addr,
+                                            sizeof(struct radv_accel_struct_serialization_header));
                instance_addr = nir_iadd(&b, dst_base_addr, nir_u2u64(&b, instance_addr));
 
                nir_build_store_global(&b, nir_channels(&b, value, 3), instance_addr,
@@ -1586,11 +1540,9 @@ build_copy_shader(struct radv_device *dev)
             }
             nir_push_else(&b, NULL);
             {
-               nir_ssa_def *instance_addr =
-                  nir_imul(&b, instance_id, nir_imm_int(&b, sizeof(uint64_t)));
-               instance_addr =
-                  nir_iadd(&b, instance_addr,
-                           nir_imm_int(&b, sizeof(struct radv_accel_struct_serialization_header)));
+               nir_ssa_def *instance_addr = nir_imul_imm(&b, instance_id, sizeof(uint64_t));
+               instance_addr = nir_iadd_imm(&b, instance_addr,
+                                            sizeof(struct radv_accel_struct_serialization_header));
                instance_addr = nir_iadd(&b, src_base_addr, nir_u2u64(&b, instance_addr));
 
                nir_ssa_def *instance_value =
diff --git a/src/amd/vulkan/radv_meta.c b/src/amd/vulkan/radv_meta.c
index 5628eedd806..861fc939e08 100644
--- a/src/amd/vulkan/radv_meta.c
+++ b/src/amd/vulkan/radv_meta.c
@@ -690,8 +690,8 @@ radv_meta_gen_rect_vertices_comp2(nir_builder *vs_b, nir_ssa_def *comp2)
    /* so channel 0 is vertex_id != 2 ? -1.0 : 1.0
       channel 1 is vertex id != 1 ? -1.0 : 1.0 */
 
-   nir_ssa_def *c0cmp = nir_ine(vs_b, vertex_id, nir_imm_int(vs_b, 2));
-   nir_ssa_def *c1cmp = nir_ine(vs_b, vertex_id, nir_imm_int(vs_b, 1));
+   nir_ssa_def *c0cmp = nir_ine_imm(vs_b, vertex_id, 2);
+   nir_ssa_def *c1cmp = nir_ine_imm(vs_b, vertex_id, 1);
 
    nir_ssa_def *comp[4];
    comp[0] = nir_bcsel(vs_b, c0cmp, nir_imm_float(vs_b, -1.0), nir_imm_float(vs_b, 1.0));
@@ -846,6 +846,6 @@ radv_break_on_count(nir_builder *b, nir_variable *var, nir_ssa_def *count)
    nir_jump(b, nir_jump_break);
    nir_pop_if(b, NULL);
 
-   counter = nir_iadd(b, counter, nir_imm_int(b, 1));
+   counter = nir_iadd_imm(b, counter, 1);
    nir_store_var(b, var, counter, 0x1);
 }
diff --git a/src/amd/vulkan/radv_meta_blit.c b/src/amd/vulkan/radv_meta_blit.c
index 0ab5124086a..2d7ab1e473c 100644
--- a/src/amd/vulkan/radv_meta_blit.c
+++ b/src/amd/vulkan/radv_meta_blit.c
@@ -64,8 +64,8 @@ build_nir_vertex_shader(struct radv_device *dev)
    /* so channel 0 is vertex_id != 2 ? src_x : src_x + w
       channel 1 is vertex id != 1 ? src_y : src_y + w */
 
-   nir_ssa_def *c0cmp = nir_ine(&b, vertex_id, nir_imm_int(&b, 2));
-   nir_ssa_def *c1cmp = nir_ine(&b, vertex_id, nir_imm_int(&b, 1));
+   nir_ssa_def *c0cmp = nir_ine_imm(&b, vertex_id, 2);
+   nir_ssa_def *c1cmp = nir_ine_imm(&b, vertex_id, 1);
 
    nir_ssa_def *comp[4];
    comp[0] = nir_bcsel(&b, c0cmp, nir_channel(&b, src_box, 0), nir_channel(&b, src_box, 2));
diff --git a/src/amd/vulkan/radv_meta_blit2d.c b/src/amd/vulkan/radv_meta_blit2d.c
index 15394078ad4..798ca9c6c3e 100644
--- a/src/amd/vulkan/radv_meta_blit2d.c
+++ b/src/amd/vulkan/radv_meta_blit2d.c
@@ -411,8 +411,8 @@ build_nir_vertex_shader(struct radv_device *device)
    /* so channel 0 is vertex_id != 2 ? src_x : src_x + w
       channel 1 is vertex id != 1 ? src_y : src_y + w */
 
-   nir_ssa_def *c0cmp = nir_ine(&b, vertex_id, nir_imm_int(&b, 2));
-   nir_ssa_def *c1cmp = nir_ine(&b, vertex_id, nir_imm_int(&b, 1));
+   nir_ssa_def *c0cmp = nir_ine_imm(&b, vertex_id, 2);
+   nir_ssa_def *c1cmp = nir_ine_imm(&b, vertex_id, 1);
 
    nir_ssa_def *comp[2];
    comp[0] = nir_bcsel(&b, c0cmp, nir_channel(&b, src_box, 0), nir_channel(&b, src_box, 2));
diff --git a/src/amd/vulkan/radv_meta_buffer.c b/src/amd/vulkan/radv_meta_buffer.c
index e6107d31ac7..fdff0368a6c 100644
--- a/src/amd/vulkan/radv_meta_buffer.c
+++ b/src/amd/vulkan/radv_meta_buffer.c
@@ -12,7 +12,7 @@ build_buffer_fill_shader(struct radv_device *dev)
 
    nir_ssa_def *global_id = get_global_ids(&b, 1);
 
-   nir_ssa_def *offset = nir_imul(&b, global_id, nir_imm_int(&b, 16));
+   nir_ssa_def *offset = nir_imul_imm(&b, global_id, 16);
    offset = nir_channel(&b, offset, 0);
 
    nir_ssa_def *dst_buf = radv_meta_load_descriptor(&b, 0, 0);
@@ -34,7 +34,7 @@ build_buffer_copy_shader(struct radv_device *dev)
 
    nir_ssa_def *global_id = get_global_ids(&b, 1);
 
-   nir_ssa_def *offset = nir_imul(&b, global_id, nir_imm_int(&b, 16));
+   nir_ssa_def *offset = nir_imul_imm(&b, global_id, 16);
    offset = nir_channel(&b, offset, 0);
 
    nir_ssa_def *dst_buf = radv_meta_load_descriptor(&b, 0, 0);
diff --git a/src/amd/vulkan/radv_meta_bufimage.c b/src/amd/vulkan/radv_meta_bufimage.c
index e6117672d18..2d0b99e7ce8 100644
--- a/src/amd/vulkan/radv_meta_bufimage.c
+++ b/src/amd/vulkan/radv_meta_bufimage.c
@@ -423,9 +423,8 @@ build_nir_btoi_r32g32b32_compute_shader(struct radv_device *dev)
 
    nir_ssa_def *img_coord = nir_iadd(&b, global_id, offset);
 
-   nir_ssa_def *global_pos =
-      nir_iadd(&b, nir_imul(&b, nir_channel(&b, img_coord, 1), pitch),
-               nir_imul(&b, nir_channel(&b, img_coord, 0), nir_imm_int(&b, 3)));
+   nir_ssa_def *global_pos = nir_iadd(&b, nir_imul(&b, nir_channel(&b, img_coord, 1), pitch),
+                                      nir_imul_imm(&b, nir_channel(&b, img_coord, 0), 3));
 
    nir_ssa_def *input_img_deref = &nir_build_deref_var(&b, input_img)->dest.ssa;
 
@@ -447,7 +446,7 @@ build_nir_btoi_r32g32b32_compute_shader(struct radv_device *dev)
    nir_ssa_def *outval = &tex->dest.ssa;
 
    for (int chan = 0; chan < 3; chan++) {
-      nir_ssa_def *local_pos = nir_iadd(&b, global_pos, nir_imm_int(&b, chan));
+      nir_ssa_def *local_pos = nir_iadd_imm(&b, global_pos, chan);
 
       nir_ssa_def *coord = nir_vec4(&b, local_pos, local_pos, local_pos, local_pos);
 
@@ -774,15 +773,15 @@ build_nir_itoi_r32g32b32_compute_shader(struct radv_device *dev)
 
    nir_ssa_def *src_global_pos =
       nir_iadd(&b, nir_imul(&b, nir_channel(&b, src_img_coord, 1), src_stride),
-               nir_imul(&b, nir_channel(&b, src_img_coord, 0), nir_imm_int(&b, 3)));
+               nir_imul_imm(&b, nir_channel(&b, src_img_coord, 0), 3));
 
    nir_ssa_def *dst_global_pos =
       nir_iadd(&b, nir_imul(&b, nir_channel(&b, dst_img_coord, 1), dst_stride),
-               nir_imul(&b, nir_channel(&b, dst_img_coord, 0), nir_imm_int(&b, 3)));
+               nir_imul_imm(&b, nir_channel(&b, dst_img_coord, 0), 3));
 
    for (int chan = 0; chan < 3; chan++) {
       /* src */
-      nir_ssa_def *src_local_pos = nir_iadd(&b, src_global_pos, nir_imm_int(&b, chan));
+      nir_ssa_def *src_local_pos = nir_iadd_imm(&b, src_global_pos, chan);
       nir_ssa_def *input_img_deref = &nir_build_deref_var(&b, input_img)->dest.ssa;
 
       nir_tex_instr *tex = nir_tex_instr_create(b.shader, 3);
@@ -803,7 +802,7 @@ build_nir_itoi_r32g32b32_compute_shader(struct radv_device *dev)
       nir_ssa_def *outval = &tex->dest.ssa;
 
       /* dst */
-      nir_ssa_def *dst_local_pos = nir_iadd(&b, dst_global_pos, nir_imm_int(&b, chan));
+      nir_ssa_def *dst_local_pos = nir_iadd_imm(&b, dst_global_pos, chan);
 
       nir_ssa_def *dst_coord =
          nir_vec4(&b, dst_local_pos, dst_local_pos, dst_local_pos, dst_local_pos);
@@ -1081,10 +1080,10 @@ build_nir_cleari_r32g32b32_compute_shader(struct radv_device *dev)
    nir_ssa_def *global_y = nir_channel(&b, global_id, 1);
 
    nir_ssa_def *global_pos =
-      nir_iadd(&b, nir_imul(&b, global_y, stride), nir_imul(&b, global_x, nir_imm_int(&b, 3)));
+      nir_iadd(&b, nir_imul(&b, global_y, stride), nir_imul_imm(&b, global_x, 3));
 
    for (unsigned chan = 0; chan < 3; chan++) {
-      nir_ssa_def *local_pos = nir_iadd(&b, global_pos, nir_imm_int(&b, chan));
+      nir_ssa_def *local_pos = nir_iadd_imm(&b, global_pos, chan);
 
       nir_ssa_def *coord = nir_vec4(&b, local_pos, local_pos, local_pos, local_pos);
 
diff --git a/src/amd/vulkan/radv_meta_clear.c b/src/amd/vulkan/radv_meta_clear.c
index bffde4b9f94..4c97c1fb7fc 100644
--- a/src/amd/vulkan/radv_meta_clear.c
+++ b/src/amd/vulkan/radv_meta_clear.c
@@ -912,7 +912,7 @@ build_clear_htile_mask_shader(struct radv_device *dev)
 
    nir_ssa_def *global_id = get_global_ids(&b, 1);
 
-   nir_ssa_def *offset = nir_imul(&b, global_id, nir_imm_int(&b, 16));
+   nir_ssa_def *offset = nir_imul_imm(&b, global_id, 16);
    offset = nir_channel(&b, offset, 0);
 
    nir_ssa_def *buf = radv_meta_load_descriptor(&b, 0, 0);
diff --git a/src/amd/vulkan/radv_meta_copy_vrs_htile.c b/src/amd/vulkan/radv_meta_copy_vrs_htile.c
index 329eaa93d5a..d4b87939fe5 100644
--- a/src/amd/vulkan/radv_meta_copy_vrs_htile.c
+++ b/src/amd/vulkan/radv_meta_copy_vrs_htile.c
@@ -52,7 +52,7 @@ build_copy_vrs_htile_shader(struct radv_device *device, struct radeon_surf *surf
    nir_ssa_def *global_id = get_global_ids(&b, 2);
 
    /* Multiply the coordinates by the HTILE block size. */
-   nir_ssa_def *coord = nir_imul(&b, global_id, nir_imm_ivec2(&b, 8, 8));
+   nir_ssa_def *coord = nir_imul_imm(&b, global_id, 8);
 
    /* Load constants. */
    nir_ssa_def *constants = nir_load_push_constant(&b, 3, 32, nir_imm_int(&b, 0), .range = 12);
@@ -99,15 +99,14 @@ build_copy_vrs_htile_shader(struct radv_device *device, struct radeon_surf *surf
     * VRS rate X = min(value >> 2, 1)
     * VRS rate Y = min(value & 3, 1)
     */
-   nir_ssa_def *x_rate = nir_ushr(&b, nir_channel(&b, &tex->dest.ssa, 0), nir_imm_int(&b, 2));
+   nir_ssa_def *x_rate = nir_ushr_imm(&b, nir_channel(&b, &tex->dest.ssa, 0), 2);
    x_rate = nir_umin(&b, x_rate, nir_imm_int(&b, 1));
 
-   nir_ssa_def *y_rate = nir_iand(&b, nir_channel(&b, &tex->dest.ssa, 0), nir_imm_int(&b, 3));
+   nir_ssa_def *y_rate = nir_iand_imm(&b, nir_channel(&b, &tex->dest.ssa, 0), 3);
    y_rate = nir_umin(&b, y_rate, nir_imm_int(&b, 1));
 
    /* Compute the final VRS rate. */
-   nir_ssa_def *vrs_rates = nir_ior(&b, nir_ishl(&b, y_rate, nir_imm_int(&b, 10)),
-                                    nir_ishl(&b, x_rate, nir_imm_int(&b, 6)));
+   nir_ssa_def *vrs_rates = nir_ior(&b, nir_ishl_imm(&b, y_rate, 10), nir_ishl_imm(&b, x_rate, 6));
 
    /* Load the HTILE buffer descriptor. */
    nir_ssa_def *htile_buf = radv_meta_load_descriptor(&b, 0, 1);
@@ -115,13 +114,13 @@ build_copy_vrs_htile_shader(struct radv_device *device, struct radeon_surf *surf
    /* Load the HTILE value if requested, otherwise use the default value. */
    nir_variable *htile_value = nir_local_variable_create(b.impl, glsl_int_type(), "htile_value");
 
-   nir_push_if(&b, nir_ieq(&b, read_htile_value, nir_imm_int(&b, 1)));
+   nir_push_if(&b, nir_ieq_imm(&b, read_htile_value, 1));
    {
       /* Load the existing HTILE 32-bit value for this 8x8 pixels area. */
       nir_ssa_def *input_value = nir_load_ssbo(&b, 1, 32, htile_buf, htile_addr);
 
       /* Clear the 4-bit VRS rates. */
-      nir_store_var(&b, htile_value, nir_iand(&b, input_value, nir_imm_int(&b, 0xfffff33f)), 0x1);
+      nir_store_var(&b, htile_value, nir_iand_imm(&b, input_value, 0xfffff33f), 0x1);
    }
    nir_push_else(&b, NULL);
    {
diff --git a/src/amd/vulkan/radv_meta_etc_decode.c b/src/amd/vulkan/radv_meta_etc_decode.c
index 6350525379f..4920d8a209b 100644
--- a/src/amd/vulkan/radv_meta_etc_decode.c
+++ b/src/amd/vulkan/radv_meta_etc_decode.c
@@ -49,12 +49,10 @@ flip_endian(nir_builder *b, nir_ssa_def *src, unsigned cnt)
       nir_ssa_def *intermediate[4];
       nir_ssa_def *chan = cnt == 1 ? src : nir_channel(b, src, i);
       for (unsigned j = 0; j < 4; ++j)
-         intermediate[j] = nir_ubfe(b, chan, nir_imm_int(b, 8 * j), nir_imm_int(b, 8));
-      v[i] = nir_ior(b,
-                     nir_ior(b, nir_ishl(b, intermediate[0], nir_imm_int(b, 24)),
-                             nir_ishl(b, intermediate[1], nir_imm_int(b, 16))),
-                     nir_ior(b, nir_ishl(b, intermediate[2], nir_imm_int(b, 8)),
-                             nir_ishl(b, intermediate[3], nir_imm_int(b, 0))));
+         intermediate[j] = nir_ubfe_imm(b, chan, 8 * j, 8);
+      v[i] = nir_ior(
+         b, nir_ior(b, nir_ishl_imm(b, intermediate[0], 24), nir_ishl_imm(b, intermediate[1], 16)),
+         nir_ior(b, nir_ishl_imm(b, intermediate[2], 8), nir_ishl_imm(b, intermediate[3], 0)));
    }
    return cnt == 1 ? v[0] : nir_vec(b, v, cnt);
 }
@@ -64,13 +62,13 @@ etc1_color_modifier_lookup(nir_builder *b, nir_ssa_def *x, nir_ssa_def *y)
 {
    const unsigned table[8][2] = {{2, 8},   {5, 17},  {9, 29},   {13, 42},
                                  {18, 60}, {24, 80}, {33, 106}, {47, 183}};
-   nir_ssa_def *upper = nir_ieq(b, y, nir_imm_int(b, 1));
+   nir_ssa_def *upper = nir_ieq_imm(b, y, 1);
    nir_ssa_def *result = NULL;
    for (unsigned i = 0; i < 8; ++i) {
       nir_ssa_def *tmp =
          nir_bcsel(b, upper, nir_imm_int(b, table[i][1]), nir_imm_int(b, table[i][0]));
       if (result)
-         result = nir_bcsel(b, nir_ieq(b, x, nir_imm_int(b, i)), tmp, result);
+         result = nir_bcsel(b, nir_ieq_imm(b, x, i), tmp, result);
       else
          result = tmp;
    }
@@ -84,7 +82,7 @@ etc2_distance_lookup(nir_builder *b, nir_ssa_def *x)
    nir_ssa_def *result = NULL;
    for (unsigned i = 0; i < 8; ++i) {
       if (result)
-         result = nir_bcsel(b, nir_ieq(b, x, nir_imm_int(b, i)), nir_imm_int(b, table[i]), result);
+         result = nir_bcsel(b, nir_ieq_imm(b, x, i), nir_imm_int(b, table[i]), result);
       else
          result = nir_imm_int(b, table[i]);
    }
@@ -100,20 +98,19 @@ etc1_alpha_modifier_lookup(nir_builder *b, nir_ssa_def *x, nir_ssa_def *y)
    for (unsigned i = 0; i < 16; ++i) {
       nir_ssa_def *tmp = nir_imm_int(b, table[i]);
       if (result)
-         result = nir_bcsel(b, nir_ieq(b, x, nir_imm_int(b, i)), tmp, result);
+         result = nir_bcsel(b, nir_ieq_imm(b, x, i), tmp, result);
       else
          result = tmp;
    }
-   return nir_ubfe(b, result, nir_imul(b, y, nir_imm_int(b, 4)), nir_imm_int(b, 4));
+   return nir_ubfe(b, result, nir_imul_imm(b, y, 4), nir_imm_int(b, 4));
 }
 
 static nir_ssa_def *
 etc_extend(nir_builder *b, nir_ssa_def *v, int bits)
 {
    if (bits == 4)
-      return nir_imul(b, v, nir_imm_int(b, 0x11));
-   return nir_ior(b, nir_ishl(b, v, nir_imm_int(b, 8 - bits)),
-                  nir_ushr(b, v, nir_imm_int(b, bits - (8 - bits))));
+      return nir_imul_imm(b, v, 0x11);
+   return nir_ior(b, nir_ishl_imm(b, v, 8 - bits), nir_ushr_imm(b, v, bits - (8 - bits)));
 }
 
 static nir_ssa_def *
@@ -123,29 +120,28 @@ decode_etc2_alpha(struct nir_builder *b, nir_ssa_def *alpha_payload, nir_ssa_def
    alpha_payload = flip_endian(b, alpha_payload, 2);
    nir_ssa_def *alpha_x = nir_channel(b, alpha_payload, 1);
    nir_ssa_def *alpha_y = nir_channel(b, alpha_payload, 0);
-   nir_ssa_def *bit_offset =
-      nir_isub(b, nir_imm_int(b, 45), nir_imul(b, nir_imm_int(b, 3), linear_pixel));
-   nir_ssa_def *base = nir_ubfe(b, alpha_y, nir_imm_int(b, 24), nir_imm_int(b, 8));
-   nir_ssa_def *multiplier = nir_ubfe(b, alpha_y, nir_imm_int(b, 20), nir_imm_int(b, 4));
-   nir_ssa_def *table = nir_ubfe(b, alpha_y, nir_imm_int(b, 16), nir_imm_int(b, 4));
+   nir_ssa_def *bit_offset = nir_isub_imm(b, 45, nir_imul_imm(b, linear_pixel, 3));
+   nir_ssa_def *base = nir_ubfe_imm(b, alpha_y, 24, 8);
+   nir_ssa_def *multiplier = nir_ubfe_imm(b, alpha_y, 20, 4);
+   nir_ssa_def *table = nir_ubfe_imm(b, alpha_y, 16, 4);
 
    if (eac) {
-      nir_ssa_def *signed_base = nir_ibfe(b, alpha_y, nir_imm_int(b, 24), nir_imm_int(b, 8));
-      signed_base = nir_imul(b, signed_base, nir_imm_int(b, 8));
-      base = nir_iadd(b, nir_imul(b, base, nir_imm_int(b, 8)), nir_imm_int(b, 4));
+      nir_ssa_def *signed_base = nir_ibfe_imm(b, alpha_y, 24, 8);
+      signed_base = nir_imul_imm(b, signed_base, 8);
+      base = nir_iadd_imm(b, nir_imul_imm(b, base, 8), 4);
       base = nir_bcsel(b, is_signed, signed_base, base);
-      multiplier = nir_imax(b, nir_imul(b, multiplier, nir_imm_int(b, 8)), nir_imm_int(b, 1));
+      multiplier = nir_imax(b, nir_imul_imm(b, multiplier, 8), nir_imm_int(b, 1));
    }
 
    nir_ssa_def *lsb_index =
       nir_ubfe(b, nir_bcsel(b, nir_uge(b, bit_offset, nir_imm_int(b, 32)), alpha_y, alpha_x),
-               nir_iand(b, bit_offset, nir_imm_int(b, 31)), nir_imm_int(b, 2));
-   bit_offset = nir_iadd(b, bit_offset, nir_imm_int(b, 2));
+               nir_iand_imm(b, bit_offset, 31), nir_imm_int(b, 2));
+   bit_offset = nir_iadd_imm(b, bit_offset, 2);
    nir_ssa_def *msb =
       nir_ubfe(b, nir_bcsel(b, nir_uge(b, bit_offset, nir_imm_int(b, 32)), alpha_y, alpha_x),
-               nir_iand(b, bit_offset, nir_imm_int(b, 31)), nir_imm_int(b, 1));
-   nir_ssa_def *mod = nir_ixor(b, etc1_alpha_modifier_lookup(b, table, lsb_index),
-                               nir_isub(b, msb, nir_imm_int(b, 1)));
+               nir_iand_imm(b, bit_offset, 31), nir_imm_int(b, 1));
+   nir_ssa_def *mod =
+      nir_ixor(b, etc1_alpha_modifier_lookup(b, table, lsb_index), nir_iadd_imm(b, msb, -1));
    nir_ssa_def *a = nir_iadd(b, base, nir_imul(b, mod, multiplier));
 
    nir_ssa_def *low_bound = nir_imm_int(b, 0);
@@ -204,7 +200,7 @@ build_shader(struct radv_device *dev)
    nir_ssa_def *offset = nir_channels(&b, consts, 7);
    nir_ssa_def *format = nir_channel(&b, consts, 3);
    nir_ssa_def *image_type = nir_channel(&b, consts2, 0);
-   nir_ssa_def *is_3d = nir_ieq(&b, image_type, nir_imm_int(&b, VK_IMAGE_TYPE_3D));
+   nir_ssa_def *is_3d = nir_ieq_imm(&b, image_type, VK_IMAGE_TYPE_3D);
    nir_ssa_def *coord = nir_iadd(&b, global_id, offset);
    nir_ssa_def *src_coord =
       nir_vec3(&b, nir_ushr_imm(&b, nir_channel(&b, coord, 0), 2),
@@ -256,10 +252,9 @@ build_shader(struct radv_device *dev)
    }
    nir_pop_if(&b, NULL);
 
-   nir_ssa_def *pixel_coord = nir_iand(&b, nir_channels(&b, coord, 3), nir_imm_ivec2(&b, 3, 3));
-   nir_ssa_def *linear_pixel =
-      nir_iadd(&b, nir_imul(&b, nir_channel(&b, pixel_coord, 0), nir_imm_int(&b, 4)),
-               nir_channel(&b, pixel_coord, 1));
+   nir_ssa_def *pixel_coord = nir_iand_imm(&b, nir_channels(&b, coord, 3), 3);
+   nir_ssa_def *linear_pixel = nir_iadd(&b, nir_imul_imm(&b, nir_channel(&b, pixel_coord, 0), 4),
+                                        nir_channel(&b, pixel_coord, 1));
 
    nir_ssa_def *payload = nir_load_var(&b, payload_var);
    nir_variable *color =
@@ -278,8 +273,7 @@ build_shader(struct radv_device *dev)
       color_payload = flip_endian(&b, color_payload, 2);
       nir_ssa_def *color_y = nir_channel(&b, color_payload, 0);
       nir_ssa_def *color_x = nir_channel(&b, color_payload, 1);
-      nir_ssa_def *flip =
-         nir_ine(&b, nir_iand(&b, color_y, nir_imm_int(&b, 1)), nir_imm_int(&b, 0));
+      nir_ssa_def *flip = nir_ine_imm(&b, nir_iand_imm(&b, color_y, 1), 0);
       nir_ssa_def *subblock = nir_ushr_imm(
          &b, nir_bcsel(&b, flip, nir_channel(&b, pixel_coord, 1), nir_channel(&b, pixel_coord, 0)),
          1);
@@ -287,8 +281,7 @@ build_shader(struct radv_device *dev)
       nir_variable *punchthrough =
          nir_variable_create(b.shader, nir_var_shader_temp, glsl_bool_type(), "punchthrough");
       nir_ssa_def *punchthrough_init =
-         nir_iand(&b, alpha_bits_1,
-                  nir_ieq(&b, nir_iand(&b, color_y, nir_imm_int(&b, 2)), nir_imm_int(&b, 0)));
+         nir_iand(&b, alpha_bits_1, nir_ieq_imm(&b, nir_iand_imm(&b, color_y, 2), 0));
       nir_store_var(&b, punchthrough, punchthrough_init, 0x1);
 
       nir_variable *etc1_compat =
@@ -317,72 +310,63 @@ build_shader(struct radv_device *dev)
       nir_store_var(&b, rgb_result, nir_imm_ivec3(&b, 255, 0, 0), 0x7);
 
       nir_ssa_def *msb =
-         nir_iand(&b, nir_ushr(&b, color_x, nir_iadd(&b, nir_imm_int(&b, 15), linear_pixel)),
-                  nir_imm_int(&b, 2));
-      nir_ssa_def *lsb = nir_iand(&b, nir_ushr(&b, color_x, linear_pixel), nir_imm_int(&b, 1));
+         nir_iand_imm(&b, nir_ushr(&b, color_x, nir_iadd_imm(&b, linear_pixel, 15)), 2);
+      nir_ssa_def *lsb = nir_iand_imm(&b, nir_ushr(&b, color_x, linear_pixel), 1);
 
-      nir_push_if(
-         &b, nir_iand(&b, nir_inot(&b, alpha_bits_1),
-                      nir_ieq(&b, nir_iand(&b, color_y, nir_imm_int(&b, 2)), nir_imm_int(&b, 0))));
+      nir_push_if(&b, nir_iand(&b, nir_inot(&b, alpha_bits_1),
+                               nir_ieq_imm(&b, nir_iand_imm(&b, color_y, 2), 0)));
       {
          nir_store_var(&b, etc1_compat, nir_imm_bool(&b, true), 1);
          nir_ssa_def *tmp[3];
          for (unsigned i = 0; i < 3; ++i)
-            tmp[i] =
-               etc_extend(&b,
-                          nir_iand(&b,
-                                   nir_ushr(&b, color_y,
-                                            nir_isub(&b, nir_imm_int(&b, 28 - 8 * i),
-                                                     nir_imul(&b, subblock, nir_imm_int(&b, 4)))),
-                                   nir_imm_int(&b, 0xf)),
-                          4);
+            tmp[i] = etc_extend(
+               &b,
+               nir_iand_imm(&b,
+                            nir_ushr(&b, color_y,
+                                     nir_isub_imm(&b, 28 - 8 * i, nir_imul_imm(&b, subblock, 4))),
+                            0xf),
+               4);
          nir_store_var(&b, base_rgb, nir_vec(&b, tmp, 3), 0x7);
       }
       nir_push_else(&b, NULL);
       {
-         nir_ssa_def *rb = nir_ubfe(&b, color_y, nir_imm_int(&b, 27), nir_imm_int(&b, 5));
-         nir_ssa_def *rd = nir_ibfe(&b, color_y, nir_imm_int(&b, 24), nir_imm_int(&b, 3));
-         nir_ssa_def *gb = nir_ubfe(&b, color_y, nir_imm_int(&b, 19), nir_imm_int(&b, 5));
-         nir_ssa_def *gd = nir_ibfe(&b, color_y, nir_imm_int(&b, 16), nir_imm_int(&b, 3));
-         nir_ssa_def *bb = nir_ubfe(&b, color_y, nir_imm_int(&b, 11), nir_imm_int(&b, 5));
-         nir_ssa_def *bd = nir_ibfe(&b, color_y, nir_imm_int(&b, 8), nir_imm_int(&b, 3));
+         nir_ssa_def *rb = nir_ubfe_imm(&b, color_y, 27, 5);
+         nir_ssa_def *rd = nir_ibfe_imm(&b, color_y, 24, 3);
+         nir_ssa_def *gb = nir_ubfe_imm(&b, color_y, 19, 5);
+         nir_ssa_def *gd = nir_ibfe_imm(&b, color_y, 16, 3);
+         nir_ssa_def *bb = nir_ubfe_imm(&b, color_y, 11, 5);
+         nir_ssa_def *bd = nir_ibfe_imm(&b, color_y, 8, 3);
          nir_ssa_def *r1 = nir_iadd(&b, rb, rd);
          nir_ssa_def *g1 = nir_iadd(&b, gb, gd);
          nir_ssa_def *b1 = nir_iadd(&b, bb, bd);
 
          nir_push_if(&b, nir_ult(&b, nir_imm_int(&b, 31), r1));
          {
-            nir_ssa_def *r0 =
-               nir_ior(&b, nir_ubfe(&b, color_y, nir_imm_int(&b, 24), nir_imm_int(&b, 2)),
-                       nir_ishl(&b, nir_ubfe(&b, color_y, nir_imm_int(&b, 27), nir_imm_int(&b, 2)),
-                                nir_imm_int(&b, 2)));
-            nir_ssa_def *g0 = nir_ubfe(&b, color_y, nir_imm_int(&b, 20), nir_imm_int(&b, 4));
-            nir_ssa_def *b0 = nir_ubfe(&b, color_y, nir_imm_int(&b, 16), nir_imm_int(&b, 4));
-            nir_ssa_def *r2 = nir_ubfe(&b, color_y, nir_imm_int(&b, 12), nir_imm_int(&b, 4));
-            nir_ssa_def *g2 = nir_ubfe(&b, color_y, nir_imm_int(&b, 8), nir_imm_int(&b, 4));
-            nir_ssa_def *b2 = nir_ubfe(&b, color_y, nir_imm_int(&b, 4), nir_imm_int(&b, 4));
-            nir_ssa_def *da =
-               nir_ior(&b,
-                       nir_ishl(&b, nir_ubfe(&b, color_y, nir_imm_int(&b, 2), nir_imm_int(&b, 2)),
-                                nir_imm_int(&b, 1)),
-                       nir_iand(&b, color_y, nir_imm_int(&b, 1)));
+            nir_ssa_def *r0 = nir_ior(&b, nir_ubfe_imm(&b, color_y, 24, 2),
+                                      nir_ishl_imm(&b, nir_ubfe_imm(&b, color_y, 27, 2), 2));
+            nir_ssa_def *g0 = nir_ubfe_imm(&b, color_y, 20, 4);
+            nir_ssa_def *b0 = nir_ubfe_imm(&b, color_y, 16, 4);
+            nir_ssa_def *r2 = nir_ubfe_imm(&b, color_y, 12, 4);
+            nir_ssa_def *g2 = nir_ubfe_imm(&b, color_y, 8, 4);
+            nir_ssa_def *b2 = nir_ubfe_imm(&b, color_y, 4, 4);
+            nir_ssa_def *da = nir_ior(&b, nir_ishl_imm(&b, nir_ubfe_imm(&b, color_y, 2, 2), 1),
+                                      nir_iand_imm(&b, color_y, 1));
             nir_ssa_def *dist = etc2_distance_lookup(&b, da);
             nir_ssa_def *index = nir_ior(&b, lsb, msb);
 
             nir_store_var(&b, punchthrough,
                           nir_iand(&b, nir_load_var(&b, punchthrough),
-                                   nir_ieq(&b, nir_iadd(&b, lsb, msb), nir_imm_int(&b, 2))),
+                                   nir_ieq_imm(&b, nir_iadd(&b, lsb, msb), 2)),
                           0x1);
-            nir_push_if(&b, nir_ieq(&b, index, nir_imm_int(&b, 0)));
+            nir_push_if(&b, nir_ieq_imm(&b, index, 0));
             {
                nir_store_var(&b, rgb_result, etc_extend(&b, nir_vec3(&b, r0, g0, b0), 4), 0x7);
             }
             nir_push_else(&b, NULL);
             {
 
-               nir_ssa_def *tmp =
-                  nir_iadd(&b, etc_extend(&b, nir_vec3(&b, r2, g2, b2), 4),
-                           nir_imul(&b, dist, nir_isub(&b, nir_imm_int(&b, 2), index)));
+               nir_ssa_def *tmp = nir_iadd(&b, etc_extend(&b, nir_vec3(&b, r2, g2, b2), 4),
+                                           nir_imul(&b, dist, nir_isub_imm(&b, 2, index)));
                nir_store_var(&b, rgb_result, tmp, 0x7);
             }
             nir_pop_if(&b, NULL);
@@ -390,64 +374,51 @@ build_shader(struct radv_device *dev)
          nir_push_else(&b, NULL);
          nir_push_if(&b, nir_ult(&b, nir_imm_int(&b, 31), g1));
          {
-            nir_ssa_def *r0 = nir_ubfe(&b, color_y, nir_imm_int(&b, 27), nir_imm_int(&b, 4));
-            nir_ssa_def *g0 = nir_ior(
-               &b,
-               nir_ishl(&b, nir_ubfe(&b, color_y, nir_imm_int(&b, 24), nir_imm_int(&b, 3)),
-                        nir_imm_int(&b, 1)),
-               nir_iand(&b, nir_ushr(&b, color_y, nir_imm_int(&b, 20)), nir_imm_int(&b, 1)));
-            nir_ssa_def *b0 = nir_ior(
-               &b, nir_ubfe(&b, color_y, nir_imm_int(&b, 15), nir_imm_int(&b, 3)),
-               nir_iand(&b, nir_ushr(&b, color_y, nir_imm_int(&b, 16)), nir_imm_int(&b, 8)));
-            nir_ssa_def *r2 = nir_ubfe(&b, color_y, nir_imm_int(&b, 11), nir_imm_int(&b, 4));
-            nir_ssa_def *g2 = nir_ubfe(&b, color_y, nir_imm_int(&b, 7), nir_imm_int(&b, 4));
-            nir_ssa_def *b2 = nir_ubfe(&b, color_y, nir_imm_int(&b, 3), nir_imm_int(&b, 4));
-            nir_ssa_def *da = nir_iand(&b, color_y, nir_imm_int(&b, 4));
-            nir_ssa_def *db = nir_iand(&b, color_y, nir_imm_int(&b, 1));
-            nir_ssa_def *d = nir_iadd(&b, da, nir_imul(&b, db, nir_imm_int(&b, 2)));
-            nir_ssa_def *d0 = nir_iadd(&b, nir_ishl(&b, r0, nir_imm_int(&b, 16)),
-                                       nir_iadd(&b, nir_ishl(&b, g0, nir_imm_int(&b, 8)), b0));
-            nir_ssa_def *d2 = nir_iadd(&b, nir_ishl(&b, r2, nir_imm_int(&b, 16)),
-                                       nir_iadd(&b, nir_ishl(&b, g2, nir_imm_int(&b, 8)), b2));
-            d = nir_bcsel(&b, nir_uge(&b, d0, d2), nir_iadd(&b, d, nir_imm_int(&b, 1)), d);
+            nir_ssa_def *r0 = nir_ubfe_imm(&b, color_y, 27, 4);
+            nir_ssa_def *g0 = nir_ior(&b, nir_ishl_imm(&b, nir_ubfe_imm(&b, color_y, 24, 3), 1),
+                                      nir_iand_imm(&b, nir_ushr_imm(&b, color_y, 20), 1));
+            nir_ssa_def *b0 = nir_ior(&b, nir_ubfe_imm(&b, color_y, 15, 3),
+                                      nir_iand_imm(&b, nir_ushr_imm(&b, color_y, 16), 8));
+            nir_ssa_def *r2 = nir_ubfe_imm(&b, color_y, 11, 4);
+            nir_ssa_def *g2 = nir_ubfe_imm(&b, color_y, 7, 4);
+            nir_ssa_def *b2 = nir_ubfe_imm(&b, color_y, 3, 4);
+            nir_ssa_def *da = nir_iand_imm(&b, color_y, 4);
+            nir_ssa_def *db = nir_iand_imm(&b, color_y, 1);
+            nir_ssa_def *d = nir_iadd(&b, da, nir_imul_imm(&b, db, 2));
+            nir_ssa_def *d0 =
+               nir_iadd(&b, nir_ishl_imm(&b, r0, 16), nir_iadd(&b, nir_ishl_imm(&b, g0, 8), b0));
+            nir_ssa_def *d2 =
+               nir_iadd(&b, nir_ishl_imm(&b, r2, 16), nir_iadd(&b, nir_ishl_imm(&b, g2, 8), b2));
+            d = nir_bcsel(&b, nir_uge(&b, d0, d2), nir_iadd_imm(&b, d, 1), d);
             nir_ssa_def *dist = etc2_distance_lookup(&b, d);
-            nir_ssa_def *base = nir_bcsel(&b, nir_ine(&b, msb, nir_imm_int(&b, 0)),
-                                          nir_vec3(&b, r2, g2, b2), nir_vec3(&b, r0, g0, b0));
+            nir_ssa_def *base = nir_bcsel(&b, nir_ine_imm(&b, msb, 0), nir_vec3(&b, r2, g2, b2),
+                                          nir_vec3(&b, r0, g0, b0));
             base = etc_extend(&b, base, 4);
-            base = nir_iadd(
-               &b, base,
-               nir_imul(&b, dist,
-                        nir_isub(&b, nir_imm_int(&b, 1), nir_imul(&b, lsb, nir_imm_int(&b, 2)))));
+            base = nir_iadd(&b, base,
+                            nir_imul(&b, dist, nir_isub_imm(&b, 1, nir_imul_imm(&b, lsb, 2))));
             nir_store_var(&b, rgb_result, base, 0x7);
             nir_store_var(&b, punchthrough,
                           nir_iand(&b, nir_load_var(&b, punchthrough),
-                                   nir_ieq(&b, nir_iadd(&b, lsb, msb), nir_imm_int(&b, 2))),
+                                   nir_ieq_imm(&b, nir_iadd(&b, lsb, msb), 2)),
                           0x1);
          }
          nir_push_else(&b, NULL);
          nir_push_if(&b, nir_ult(&b, nir_imm_int(&b, 31), b1));
          {
-            nir_ssa_def *r0 = nir_ubfe(&b, color_y, nir_imm_int(&b, 25), nir_imm_int(&b, 6));
-            nir_ssa_def *g0 = nir_ior(
-               &b, nir_ubfe(&b, color_y, nir_imm_int(&b, 17), nir_imm_int(&b, 6)),
-               nir_iand(&b, nir_ushr(&b, color_y, nir_imm_int(&b, 18)), nir_imm_int(&b, 0x40)));
-            nir_ssa_def *b0 = nir_ior(
-               &b,
-               nir_ishl(&b, nir_ubfe(&b, color_y, nir_imm_int(&b, 11), nir_imm_int(&b, 2)),
-                        nir_imm_int(&b, 3)),
-               nir_ior(
-                  &b,
-                  nir_iand(&b, nir_ushr(&b, color_y, nir_imm_int(&b, 11)), nir_imm_int(&b, 0x20)),
-                  nir_ubfe(&b, color_y, nir_imm_int(&b, 7), nir_imm_int(&b, 3))));
-            nir_ssa_def *rh =
-               nir_ior(&b, nir_iand(&b, color_y, nir_imm_int(&b, 1)),
-                       nir_ishl(&b, nir_ubfe(&b, color_y, nir_imm_int(&b, 2), nir_imm_int(&b, 5)),
-                                nir_imm_int(&b, 1)));
-            nir_ssa_def *rv = nir_ubfe(&b, color_x, nir_imm_int(&b, 13), nir_imm_int(&b, 6));
-            nir_ssa_def *gh = nir_ubfe(&b, color_x, nir_imm_int(&b, 25), nir_imm_int(&b, 7));
-            nir_ssa_def *gv = nir_ubfe(&b, color_x, nir_imm_int(&b, 6), nir_imm_int(&b, 7));
-            nir_ssa_def *bh = nir_ubfe(&b, color_x, nir_imm_int(&b, 19), nir_imm_int(&b, 6));
-            nir_ssa_def *bv = nir_ubfe(&b, color_x, nir_imm_int(&b, 0), nir_imm_int(&b, 6));
+            nir_ssa_def *r0 = nir_ubfe_imm(&b, color_y, 25, 6);
+            nir_ssa_def *g0 = nir_ior(&b, nir_ubfe_imm(&b, color_y, 17, 6),
+                                      nir_iand_imm(&b, nir_ushr_imm(&b, color_y, 18), 0x40));
+            nir_ssa_def *b0 =
+               nir_ior(&b, nir_ishl_imm(&b, nir_ubfe_imm(&b, color_y, 11, 2), 3),
+                       nir_ior(&b, nir_iand_imm(&b, nir_ushr_imm(&b, color_y, 11), 0x20),
+                               nir_ubfe_imm(&b, color_y, 7, 3)));
+            nir_ssa_def *rh = nir_ior(&b, nir_iand_imm(&b, color_y, 1),
+                                      nir_ishl_imm(&b, nir_ubfe_imm(&b, color_y, 2, 5), 1));
+            nir_ssa_def *rv = nir_ubfe_imm(&b, color_x, 13, 6);
+            nir_ssa_def *gh = nir_ubfe_imm(&b, color_x, 25, 7);
+            nir_ssa_def *gv = nir_ubfe_imm(&b, color_x, 6, 7);
+            nir_ssa_def *bh = nir_ubfe_imm(&b, color_x, 19, 6);
+            nir_ssa_def *bv = nir_ubfe_imm(&b, color_x, 0, 6);
 
             r0 = etc_extend(&b, r0, 6);
             g0 = etc_extend(&b, g0, 7);
@@ -464,16 +435,14 @@ build_shader(struct radv_device *dev)
                                        nir_channel(&b, pixel_coord, 0));
             nir_ssa_def *dy = nir_imul(&b, nir_isub(&b, nir_vec3(&b, rv, gv, bv), rgb),
                                        nir_channel(&b, pixel_coord, 1));
-            rgb = nir_iadd(&b, rgb,
-                           nir_ishr(&b, nir_iadd(&b, nir_iadd(&b, dx, dy), nir_imm_int(&b, 2)),
-                                    nir_imm_int(&b, 2)));
+            rgb = nir_iadd(&b, rgb, nir_ishr_imm(&b, nir_iadd_imm(&b, nir_iadd(&b, dx, dy), 2), 2));
             nir_store_var(&b, rgb_result, rgb, 0x7);
             nir_store_var(&b, punchthrough, nir_imm_bool(&b, false), 0x1);
          }
          nir_push_else(&b, NULL);
          {
             nir_store_var(&b, etc1_compat, nir_imm_bool(&b, true), 1);
-            nir_ssa_def *subblock_b = nir_ine(&b, subblock, nir_imm_int(&b, 0));
+            nir_ssa_def *subblock_b = nir_ine_imm(&b, subblock, 0);
             nir_ssa_def *tmp[] = {
                nir_bcsel(&b, subblock_b, r1, rb),
                nir_bcsel(&b, subblock_b, g1, gb),
@@ -488,15 +457,13 @@ build_shader(struct radv_device *dev)
       nir_pop_if(&b, NULL);
       nir_push_if(&b, nir_load_var(&b, etc1_compat));
       {
-         nir_ssa_def *etc1_table_index =
-            nir_ubfe(&b, color_y,
-                     nir_isub(&b, nir_imm_int(&b, 5), nir_imul(&b, nir_imm_int(&b, 3), subblock)),
-                     nir_imm_int(&b, 3));
-         nir_ssa_def *sgn = nir_isub(&b, nir_imm_int(&b, 1), msb);
+         nir_ssa_def *etc1_table_index = nir_ubfe(
+            &b, color_y, nir_isub_imm(&b, 5, nir_imul_imm(&b, subblock, 3)), nir_imm_int(&b, 3));
+         nir_ssa_def *sgn = nir_isub_imm(&b, 1, msb);
          sgn = nir_bcsel(&b, nir_load_var(&b, punchthrough), nir_imul(&b, sgn, lsb), sgn);
          nir_store_var(&b, punchthrough,
                        nir_iand(&b, nir_load_var(&b, punchthrough),
-                                nir_ieq(&b, nir_iadd(&b, lsb, msb), nir_imm_int(&b, 2))),
+                                nir_ieq_imm(&b, nir_iadd(&b, lsb, msb), 2)),
                        0x1);
          nir_ssa_def *off =
             nir_imul(&b, etc1_color_modifier_lookup(&b, etc1_table_index, lsb), sgn);
@@ -519,9 +486,8 @@ build_shader(struct radv_device *dev)
    }
    nir_push_else(&b, NULL);
    { /* EAC */
-      nir_ssa_def *is_signed =
-         nir_ior(&b, nir_ieq(&b, format, nir_imm_int(&b, VK_FORMAT_EAC_R11_SNORM_BLOCK)),
-                 nir_ieq(&b, format, nir_imm_int(&b, VK_FORMAT_EAC_R11G11_SNORM_BLOCK)));
+      nir_ssa_def *is_signed = nir_ior(&b, nir_ieq_imm(&b, format, VK_FORMAT_EAC_R11_SNORM_BLOCK),
+                                       nir_ieq_imm(&b, format, VK_FORMAT_EAC_R11G11_SNORM_BLOCK));
       nir_ssa_def *val[4];
       for (int i = 0; i < 2; ++i) {
          val[i] = decode_etc2_alpha(&b, nir_channels(&b, payload, 3 << (2 * i)), linear_pixel, true,
diff --git a/src/amd/vulkan/radv_nir_lower_ray_queries.c b/src/amd/vulkan/radv_nir_lower_ray_queries.c
index bc7b18b69a1..574ea78ba9e 100644
--- a/src/amd/vulkan/radv_nir_lower_ray_queries.c
+++ b/src/amd/vulkan/radv_nir_lower_ray_queries.c
@@ -298,11 +298,9 @@ static void
 insert_terminate_on_first_hit(nir_builder *b, nir_ssa_def *index, struct ray_query_vars *vars,
                               bool break_on_terminate)
 {
-   nir_ssa_def *terminate_on_first_hit =
-      nir_ine(b,
-              nir_iand(b, rq_load_var(b, index, vars->flags),
-                       nir_imm_int(b, SpvRayFlagsTerminateOnFirstHitKHRMask)),
-              nir_imm_int(b, 0));
+   nir_ssa_def *terminate_on_first_hit = nir_ine_imm(
+      b, nir_iand_imm(b, rq_load_var(b, index, vars->flags), SpvRayFlagsTerminateOnFirstHitKHRMask),
+      0);
    nir_push_if(b, terminate_on_first_hit);
    {
       rq_store_var(b, index, vars->incomplete, nir_imm_bool(b, false), 0x1);
@@ -346,8 +344,7 @@ lower_rq_initialize(nir_builder *b, nir_ssa_def *index, nir_intrinsic_instr *ins
 {
    rq_store_var(b, index, vars->accel_struct, instr->src[1].ssa, 0x1);
    rq_store_var(b, index, vars->flags, instr->src[2].ssa, 0x1);
-   rq_store_var(b, index, vars->cull_mask, nir_iand(b, instr->src[3].ssa, nir_imm_int(b, 0xff)),
-                0x1);
+   rq_store_var(b, index, vars->cull_mask, nir_iand_imm(b, instr->src[3].ssa, 0xff), 0x1);
 
    rq_store_var(b, index, vars->origin, instr->src[4].ssa, 0x7);
    rq_store_var(b, index, vars->trav.origin, instr->src[4].ssa, 0x7);
@@ -366,7 +363,7 @@ lower_rq_initialize(nir_builder *b, nir_ssa_def *index, nir_intrinsic_instr *ins
 
    nir_ssa_def *accel_struct = rq_load_var(b, index, vars->accel_struct);
 
-   nir_push_if(b, nir_ine(b, accel_struct, nir_imm_int64(b, 0)));
+   nir_push_if(b, nir_ine_imm(b, accel_struct, 0));
    {
       rq_store_var(b, index, vars->trav.bvh_base, build_addr_to_node(b, accel_struct), 1);
 
@@ -400,32 +397,32 @@ lower_rq_load(nir_builder *b, nir_ssa_def *index, struct ray_query_vars *vars,
                        rq_load_var(b, index, vars->candidate.barycentrics));
    case nir_ray_query_value_intersection_candidate_aabb_opaque:
       return nir_iand(b, rq_load_var(b, index, vars->candidate.opaque),
-                      nir_ieq(b, rq_load_var(b, index, vars->candidate.intersection_type),
-                              nir_imm_int(b, intersection_type_aabb)));
+                      nir_ieq_imm(b, rq_load_var(b, index, vars->candidate.intersection_type),
+                                  intersection_type_aabb));
    case nir_ray_query_value_intersection_front_face:
       return nir_bcsel(b, committed, rq_load_var(b, index, vars->closest.frontface),
                        rq_load_var(b, index, vars->candidate.frontface));
    case nir_ray_query_value_intersection_geometry_index:
-      return nir_iand(
+      return nir_iand_imm(
          b,
          nir_bcsel(b, committed, rq_load_var(b, index, vars->closest.geometry_id_and_flags),
                    rq_load_var(b, index, vars->candidate.geometry_id_and_flags)),
-         nir_imm_int(b, 0xFFFFFF));
+         0xFFFFFF);
    case nir_ray_query_value_intersection_instance_custom_index:
-      return nir_iand(
+      return nir_iand_imm(
          b,
          nir_bcsel(b, committed, rq_load_var(b, index, vars->closest.custom_instance_and_mask),
                    rq_load_var(b, index, vars->candidate.custom_instance_and_mask)),
-         nir_imm_int(b, 0xFFFFFF));
+         0xFFFFFF);
    case nir_ray_query_value_intersection_instance_id:
       return nir_bcsel(b, committed, rq_load_var(b, index, vars->closest.instance_id),
                        rq_load_var(b, index, vars->candidate.instance_id));
    case nir_ray_query_value_intersection_instance_sbt_index:
-      return nir_iand(
+      return nir_iand_imm(
          b,
          nir_bcsel(b, committed, rq_load_var(b, index, vars->closest.sbt_offset_and_flags),
                    rq_load_var(b, index, vars->candidate.sbt_offset_and_flags)),
-         nir_imm_int(b, 0xFFFFFF));
+         0xFFFFFF);
    case nir_ray_query_value_intersection_object_ray_direction: {
       nir_ssa_def *instance_node_addr =
          nir_bcsel(b, committed, rq_load_var(b, index, vars->closest.instance_addr),
@@ -439,12 +436,12 @@ lower_rq_load(nir_builder *b, nir_ssa_def *index, struct ray_query_vars *vars,
          nir_bcsel(b, committed, rq_load_var(b, index, vars->closest.instance_addr),
                    rq_load_var(b, index, vars->candidate.instance_addr));
       nir_ssa_def *wto_matrix[] = {
-         nir_build_load_global(b, 4, 32, nir_iadd(b, instance_node_addr, nir_imm_int64(b, 16)),
-                               .align_mul = 64, .align_offset = 16),
-         nir_build_load_global(b, 4, 32, nir_iadd(b, instance_node_addr, nir_imm_int64(b, 32)),
-                               .align_mul = 64, .align_offset = 32),
-         nir_build_load_global(b, 4, 32, nir_iadd(b, instance_node_addr, nir_imm_int64(b, 48)),
-                               .align_mul = 64, .align_offset = 48)};
+         nir_build_load_global(b, 4, 32, nir_iadd_imm(b, instance_node_addr, 16), .align_mul = 64,
+                               .align_offset = 16),
+         nir_build_load_global(b, 4, 32, nir_iadd_imm(b, instance_node_addr, 32), .align_mul = 64,
+                               .align_offset = 32),
+         nir_build_load_global(b, 4, 32, nir_iadd_imm(b, instance_node_addr, 48), .align_mul = 64,
+                               .align_offset = 48)};
       return nir_build_vec3_mat_mult_pre(b, rq_load_var(b, index, vars->origin), wto_matrix);
    }
    case nir_ray_query_value_intersection_object_to_world: {
@@ -463,8 +460,7 @@ lower_rq_load(nir_builder *b, nir_ssa_def *index, struct ray_query_vars *vars,
          return nir_vec(b, vals, 3);
       }
 
-      return nir_build_load_global(
-         b, 3, 32, nir_iadd(b, instance_node_addr, nir_imm_int64(b, 92 + column * 12)));
+      return nir_build_load_global(b, 3, 32, nir_iadd_imm(b, instance_node_addr, 92 + column * 12));
    }
    case nir_ray_query_value_intersection_primitive_index:
       return nir_bcsel(b, committed, rq_load_var(b, index, vars->closest.primitive_id),
@@ -475,7 +471,7 @@ lower_rq_load(nir_builder *b, nir_ssa_def *index, struct ray_query_vars *vars,
    case nir_ray_query_value_intersection_type:
       return nir_bcsel(
          b, committed, rq_load_var(b, index, vars->closest.intersection_type),
-         nir_isub(b, rq_load_var(b, index, vars->candidate.intersection_type), nir_imm_int(b, 1)));
+         nir_iadd_imm(b, rq_load_var(b, index, vars->candidate.intersection_type), -1));
    case nir_ray_query_value_intersection_world_to_object: {
       nir_ssa_def *instance_node_addr =
          nir_bcsel(b, committed, rq_load_var(b, index, vars->closest.instance_addr),
@@ -515,44 +511,41 @@ insert_traversal_triangle_case(struct radv_device *device, nir_builder *b, nir_s
    nir_ssa_def *div = nir_vector_extract(b, result, nir_imm_int(b, 1));
    dist = nir_fdiv(b, dist, div);
    nir_ssa_def *frontface = nir_flt(b, nir_imm_float(b, 0), div);
-   nir_ssa_def *switch_ccw = nir_ine(
-      b,
-      nir_iand(b, rq_load_var(b, index, vars->candidate.sbt_offset_and_flags),
-               nir_imm_int(b, VK_GEOMETRY_INSTANCE_TRIANGLE_FRONT_COUNTERCLOCKWISE_BIT_KHR << 24)),
-      nir_imm_int(b, 0));
+   nir_ssa_def *switch_ccw =
+      nir_ine_imm(b,
+                  nir_iand_imm(b, rq_load_var(b, index, vars->candidate.sbt_offset_and_flags),
+                               VK_GEOMETRY_INSTANCE_TRIANGLE_FRONT_COUNTERCLOCKWISE_BIT_KHR << 24),
+                  0);
    frontface = nir_ixor(b, frontface, switch_ccw);
    rq_store_var(b, index, vars->candidate.frontface, frontface, 0x1);
 
-   nir_ssa_def *not_cull = nir_ieq(b,
-                                   nir_iand(b, rq_load_var(b, index, vars->flags),
-                                            nir_imm_int(b, SpvRayFlagsSkipTrianglesKHRMask)),
-                                   nir_imm_int(b, 0));
-   nir_ssa_def *not_facing_cull = nir_ieq(
+   nir_ssa_def *not_cull = nir_ieq_imm(
+      b, nir_iand_imm(b, rq_load_var(b, index, vars->flags), SpvRayFlagsSkipTrianglesKHRMask), 0);
+   nir_ssa_def *not_facing_cull = nir_ieq_imm(
       b,
       nir_iand(b, rq_load_var(b, index, vars->flags),
                nir_bcsel(b, frontface, nir_imm_int(b, SpvRayFlagsCullFrontFacingTrianglesKHRMask),
                          nir_imm_int(b, SpvRayFlagsCullBackFacingTrianglesKHRMask))),
-      nir_imm_int(b, 0));
+      0);
 
    not_cull = nir_iand(
       b, not_cull,
       nir_ior(
          b, not_facing_cull,
-         nir_ine(b,
-                 nir_iand(b, rq_load_var(b, index, vars->candidate.sbt_offset_and_flags),
-                          nir_imm_int(
-                             b, VK_GEOMETRY_INSTANCE_TRIANGLE_FACING_CULL_DISABLE_BIT_KHR << 24)),
-                 nir_imm_int(b, 0))));
+         nir_ine_imm(b,
+                     nir_iand_imm(b, rq_load_var(b, index, vars->candidate.sbt_offset_and_flags),
+                                  VK_GEOMETRY_INSTANCE_TRIANGLE_FACING_CULL_DISABLE_BIT_KHR << 24),
+                     0)));
 
    nir_push_if(b, nir_iand(b,
                            nir_iand(b, nir_flt(b, dist, rq_load_var(b, index, vars->closest.t)),
                                     nir_fge(b, dist, rq_load_var(b, index, vars->tmin))),
                            not_cull));
    {
-      nir_ssa_def *triangle_info = nir_build_load_global(
-         b, 2, 32,
-         nir_iadd(b, build_node_to_addr(device, b, bvh_node),
-                  nir_imm_int64(b, offsetof(struct radv_bvh_triangle_node, triangle_id))));
+      nir_ssa_def *triangle_info =
+         nir_build_load_global(b, 2, 32,
+                               nir_iadd_imm(b, build_node_to_addr(device, b, bvh_node),
+                                            offsetof(struct radv_bvh_triangle_node, triangle_id)));
       nir_ssa_def *primitive_id = nir_channel(b, triangle_info, 0);
       nir_ssa_def *geometry_id_and_flags = nir_channel(b, triangle_info, 1);
       nir_ssa_def *is_opaque =
@@ -560,11 +553,11 @@ insert_traversal_triangle_case(struct radv_device *device, nir_builder *b, nir_s
                        rq_load_var(b, index, vars->flags), geometry_id_and_flags);
 
       not_cull =
-         nir_ieq(b,
-                 nir_iand(b, rq_load_var(b, index, vars->flags),
-                          nir_bcsel(b, is_opaque, nir_imm_int(b, SpvRayFlagsCullOpaqueKHRMask),
-                                    nir_imm_int(b, SpvRayFlagsCullNoOpaqueKHRMask))),
-                 nir_imm_int(b, 0));
+         nir_ieq_imm(b,
+                     nir_iand(b, rq_load_var(b, index, vars->flags),
+                              nir_bcsel(b, is_opaque, nir_imm_int(b, SpvRayFlagsCullOpaqueKHRMask),
+                                        nir_imm_int(b, SpvRayFlagsCullNoOpaqueKHRMask))),
+                     0);
       nir_push_if(b, not_cull);
       {
          nir_ssa_def *divs[2] = {div, div};
@@ -599,35 +592,30 @@ insert_traversal_aabb_case(struct radv_device *device, nir_builder *b, nir_ssa_d
                            struct ray_query_vars *vars, nir_ssa_def *bvh_node)
 {
    nir_ssa_def *node_addr = build_node_to_addr(device, b, bvh_node);
-   nir_ssa_def *triangle_info =
-      nir_build_load_global(b, 2, 32, nir_iadd(b, node_addr, nir_imm_int64(b, 24)));
+   nir_ssa_def *triangle_info = nir_build_load_global(b, 2, 32, nir_iadd_imm(b, node_addr, 24));
    nir_ssa_def *primitive_id = nir_channel(b, triangle_info, 0);
    nir_ssa_def *geometry_id_and_flags = nir_channel(b, triangle_info, 1);
    nir_ssa_def *is_opaque =
       hit_is_opaque(b, rq_load_var(b, index, vars->candidate.sbt_offset_and_flags),
                     rq_load_var(b, index, vars->flags), geometry_id_and_flags);
 
-   nir_ssa_def *not_skip_aabb = nir_ieq(
-      b,
-      nir_iand(b, rq_load_var(b, index, vars->flags), nir_imm_int(b, SpvRayFlagsSkipAABBsKHRMask)),
-      nir_imm_int(b, 0));
+   nir_ssa_def *not_skip_aabb = nir_ieq_imm(
+      b, nir_iand_imm(b, rq_load_var(b, index, vars->flags), SpvRayFlagsSkipAABBsKHRMask), 0);
    nir_ssa_def *not_cull = nir_iand(
       b, not_skip_aabb,
-      nir_ieq(b,
-              nir_iand(b, rq_load_var(b, index, vars->flags),
-                       nir_bcsel(b, is_opaque, nir_imm_int(b, SpvRayFlagsCullOpaqueKHRMask),
-                                 nir_imm_int(b, SpvRayFlagsCullNoOpaqueKHRMask))),
-              nir_imm_int(b, 0)));
+      nir_ieq_imm(b,
+                  nir_iand(b, rq_load_var(b, index, vars->flags),
+                           nir_bcsel(b, is_opaque, nir_imm_int(b, SpvRayFlagsCullOpaqueKHRMask),
+                                     nir_imm_int(b, SpvRayFlagsCullNoOpaqueKHRMask))),
+                  0));
    nir_push_if(b, not_cull);
    {
       nir_ssa_def *vec3_zero = nir_channels(b, nir_imm_vec4(b, 0, 0, 0, 0), 0x7);
       nir_ssa_def *vec3_inf =
          nir_channels(b, nir_imm_vec4(b, INFINITY, INFINITY, INFINITY, 0), 0x7);
 
-      nir_ssa_def *bvh_lo =
-         nir_build_load_global(b, 3, 32, nir_iadd(b, node_addr, nir_imm_int64(b, 0)));
-      nir_ssa_def *bvh_hi =
-         nir_build_load_global(b, 3, 32, nir_iadd(b, node_addr, nir_imm_int64(b, 12)));
+      nir_ssa_def *bvh_lo = nir_build_load_global(b, 3, 32, nir_iadd_imm(b, node_addr, 0));
+      nir_ssa_def *bvh_hi = nir_build_load_global(b, 3, 32, nir_iadd_imm(b, node_addr, 12));
 
       bvh_lo = nir_fsub(b, bvh_lo, rq_load_var(b, index, vars->trav.origin));
       bvh_hi = nir_fsub(b, bvh_hi, rq_load_var(b, index, vars->trav.origin));
@@ -685,7 +673,7 @@ lower_rq_proceed(nir_builder *b, nir_ssa_def *index, struct ray_query_vars *vars
          nir_push_if(b, nir_uge(b, rq_load_var(b, index, vars->trav.top_stack),
                                 rq_load_var(b, index, vars->trav.stack)));
          {
-            nir_push_if(b, nir_ieq(b, rq_load_var(b, index, vars->trav.stack), nir_imm_int(b, 0)));
+            nir_push_if(b, nir_ieq_imm(b, rq_load_var(b, index, vars->trav.stack), 0));
             {
                rq_store_var(b, index, vars->incomplete, nir_imm_bool(b, false), 0x1);
                nir_jump(b, nir_jump_break);
@@ -703,7 +691,7 @@ lower_rq_proceed(nir_builder *b, nir_ssa_def *index, struct ray_query_vars *vars
          nir_pop_if(b, NULL);
 
          rq_store_var(b, index, vars->trav.stack,
-                      nir_isub(b, rq_load_var(b, index, vars->trav.stack), nir_imm_int(b, 1)), 1);
+                      nir_iadd_imm(b, rq_load_var(b, index, vars->trav.stack), 1), 1);
 
          nir_ssa_def *bvh_node =
             rq_load_array(b, index, vars->stack, rq_load_var(b, index, vars->trav.stack));
@@ -721,16 +709,13 @@ lower_rq_proceed(nir_builder *b, nir_ssa_def *index, struct ray_query_vars *vars
          }
 
          /* if (node.type_flags & aabb) */
-         nir_push_if(b,
-                     nir_ine(b, nir_iand(b, bvh_node_type, nir_imm_int(b, 4)), nir_imm_int(b, 0)));
+         nir_push_if(b, nir_ine_imm(b, nir_iand_imm(b, bvh_node_type, 4), 0));
          {
             /* if (node.type_flags & leaf) */
-            nir_push_if(
-               b, nir_ine(b, nir_iand(b, bvh_node_type, nir_imm_int(b, 2)), nir_imm_int(b, 0)));
+            nir_push_if(b, nir_ine_imm(b, nir_iand_imm(b, bvh_node_type, 2), 0));
             {
                /* custom */
-               nir_push_if(
-                  b, nir_ine(b, nir_iand(b, bvh_node_type, nir_imm_int(b, 1)), nir_imm_int(b, 0)));
+               nir_push_if(b, nir_ine_imm(b, nir_iand_imm(b, bvh_node_type, 1), 0));
                {
                   insert_traversal_aabb_case(device, b, index, vars, bvh_node);
                }
@@ -741,30 +726,26 @@ lower_rq_proceed(nir_builder *b, nir_ssa_def *index, struct ray_query_vars *vars
                   nir_ssa_def *instance_data = nir_build_load_global(
                      b, 4, 32, instance_node_addr, .align_mul = 64, .align_offset = 0);
                   nir_ssa_def *instance_and_mask = nir_channel(b, instance_data, 2);
-                  nir_ssa_def *instance_mask = nir_ushr(b, instance_and_mask, nir_imm_int(b, 24));
+                  nir_ssa_def *instance_mask = nir_ushr_imm(b, instance_and_mask, 24);
 
                   nir_push_if(
                      b,
-                     nir_ieq(b, nir_iand(b, instance_mask, rq_load_var(b, index, vars->cull_mask)),
-                             nir_imm_int(b, 0)));
+                     nir_ieq_imm(
+                        b, nir_iand(b, instance_mask, rq_load_var(b, index, vars->cull_mask)), 0));
                   {
                      nir_jump(b, nir_jump_continue);
                   }
                   nir_pop_if(b, NULL);
 
                   nir_ssa_def *wto_matrix[] = {
-                     nir_build_load_global(b, 4, 32,
-                                           nir_iadd(b, instance_node_addr, nir_imm_int64(b, 16)),
+                     nir_build_load_global(b, 4, 32, nir_iadd_imm(b, instance_node_addr, 16),
                                            .align_mul = 64, .align_offset = 16),
-                     nir_build_load_global(b, 4, 32,
-                                           nir_iadd(b, instance_node_addr, nir_imm_int64(b, 32)),
+                     nir_build_load_global(b, 4, 32, nir_iadd_imm(b, instance_node_addr, 32),
                                            .align_mul = 64, .align_offset = 32),
-                     nir_build_load_global(b, 4, 32,
-                                           nir_iadd(b, instance_node_addr, nir_imm_int64(b, 48)),
+                     nir_build_load_global(b, 4, 32, nir_iadd_imm(b, instance_node_addr, 48),
                                            .align_mul = 64, .align_offset = 48)};
-                  nir_ssa_def *instance_id = nir_build_load_global(
-                     b, 1, 32, nir_iadd(b, instance_node_addr, nir_imm_int64(b, 88)),
-                     .align_mul = 4, .align_offset = 0);
+                  nir_ssa_def *instance_id =
+                     nir_build_load_global(b, 1, 32, nir_iadd_imm(b, instance_node_addr, 88));
 
                   rq_store_var(b, index, vars->trav.top_stack,
                                rq_load_var(b, index, vars->trav.stack), 1);
@@ -774,11 +755,9 @@ lower_rq_proceed(nir_builder *b, nir_ssa_def *index, struct ray_query_vars *vars
                                1);
 
                   rq_store_array(b, index, vars->stack, rq_load_var(b, index, vars->trav.stack),
-                                 nir_iand(b, nir_channel(b, instance_data, 0), nir_imm_int(b, 63)),
-                                 0x1);
-                  rq_store_var(
-                     b, index, vars->trav.stack,
-                     nir_iadd(b, rq_load_var(b, index, vars->trav.stack), nir_imm_int(b, 1)), 1);
+                                 nir_iand_imm(b, nir_channel(b, instance_data, 0), 63), 0x1);
+                  rq_store_var(b, index, vars->trav.stack,
+                               nir_iadd_imm(b, rq_load_var(b, index, vars->trav.stack), 1), 1);
 
                   rq_store_var(b, index, vars->trav.origin,
                                nir_build_vec3_mat_mult_pre(b, rq_load_var(b, index, vars->origin),
@@ -817,13 +796,12 @@ lower_rq_proceed(nir_builder *b, nir_ssa_def *index, struct ray_query_vars *vars
                /* box */
                for (unsigned i = 4; i-- > 0;) {
                   nir_ssa_def *new_node = nir_vector_extract(b, result, nir_imm_int(b, i));
-                  nir_push_if(b, nir_ine(b, new_node, nir_imm_int(b, 0xffffffff)));
+                  nir_push_if(b, nir_ine_imm(b, new_node, 0xffffffff));
                   {
                      rq_store_array(b, index, vars->stack, rq_load_var(b, index, vars->trav.stack),
                                     new_node, 0x1);
-                     rq_store_var(
-                        b, index, vars->trav.stack,
-                        nir_iadd(b, rq_load_var(b, index, vars->trav.stack), nir_imm_int(b, 1)), 1);
+                     rq_store_var(b, index, vars->trav.stack,
+                                  nir_iadd_imm(b, rq_load_var(b, index, vars->trav.stack), 1), 1);
                   }
                   nir_pop_if(b, NULL);
                }
diff --git a/src/amd/vulkan/radv_pipeline.c b/src/amd/vulkan/radv_pipeline.c
index e633962cafb..7e15e763305 100644
--- a/src/amd/vulkan/radv_pipeline.c
+++ b/src/amd/vulkan/radv_pipeline.c
@@ -3824,7 +3824,7 @@ radv_adjust_vertex_fetch_alpha(nir_builder *b,
     */
    unsigned offset = alpha_adjust == ALPHA_ADJUST_SNORM ? 23u : 0u;
 
-   alpha = nir_ibfe(b, alpha, nir_imm_int(b, offset), nir_imm_int(b, 2u));
+   alpha = nir_ibfe_imm(b, alpha, offset, 2u);
 
    /* Convert back to the right type. */
    if (alpha_adjust == ALPHA_ADJUST_SNORM) {
diff --git a/src/amd/vulkan/radv_pipeline_rt.c b/src/amd/vulkan/radv_pipeline_rt.c
index b43b4e2c363..bcb54a51bf3 100644
--- a/src/amd/vulkan/radv_pipeline_rt.c
+++ b/src/amd/vulkan/radv_pipeline_rt.c
@@ -314,8 +314,7 @@ const uint32_t RADV_HIT_ATTRIB_OFFSET = -(16 + RADV_MAX_HIT_ATTRIB_SIZE);
 static void
 insert_rt_return(nir_builder *b, const struct rt_variables *vars)
 {
-   nir_store_var(b, vars->stack_ptr,
-                 nir_iadd(b, nir_load_var(b, vars->stack_ptr), nir_imm_int(b, -16)), 1);
+   nir_store_var(b, vars->stack_ptr, nir_iadd_imm(b, nir_load_var(b, vars->stack_ptr), -16), 1);
    nir_store_var(b, vars->idx,
                  nir_load_scratch(b, 1, 32, nir_load_var(b, vars->stack_ptr), .align_mul = 16), 1);
 }
@@ -346,14 +345,12 @@ load_sbt_entry(nir_builder *b, const struct rt_variables *vars, nir_ssa_def *idx
 {
    nir_ssa_def *addr = get_sbt_ptr(b, idx, binding);
 
-   nir_ssa_def *load_addr = addr;
-   if (offset)
-      load_addr = nir_iadd(b, load_addr, nir_imm_int64(b, offset));
+   nir_ssa_def *load_addr = nir_iadd_imm(b, addr, offset);
    nir_ssa_def *v_idx = nir_build_load_global(b, 1, 32, load_addr);
 
    nir_store_var(b, vars->idx, v_idx, 1);
 
-   nir_ssa_def *record_addr = nir_iadd(b, addr, nir_imm_int64(b, RADV_RT_HANDLE_SIZE));
+   nir_ssa_def *record_addr = nir_iadd_imm(b, addr, RADV_RT_HANDLE_SIZE);
    nir_store_var(b, vars->shader_record_ptr, record_addr, 1);
 }
 
@@ -376,22 +373,19 @@ lower_rt_instructions(nir_shader *shader, struct rt_variables *vars, unsigned ca
                uint32_t ret = call_idx_base + nir_intrinsic_call_idx(intr) + 1;
                b_shader.cursor = nir_instr_remove(instr);
 
-               nir_store_var(&b_shader, vars->stack_ptr,
-                             nir_iadd(&b_shader, nir_load_var(&b_shader, vars->stack_ptr),
-                                      nir_imm_int(&b_shader, size)),
-                             1);
+               nir_store_var(
+                  &b_shader, vars->stack_ptr,
+                  nir_iadd_imm(&b_shader, nir_load_var(&b_shader, vars->stack_ptr), size), 1);
                nir_store_scratch(&b_shader, nir_imm_int(&b_shader, ret),
                                  nir_load_var(&b_shader, vars->stack_ptr), .align_mul = 16);
 
                nir_store_var(&b_shader, vars->stack_ptr,
-                             nir_iadd(&b_shader, nir_load_var(&b_shader, vars->stack_ptr),
-                                      nir_imm_int(&b_shader, 16)),
+                             nir_iadd_imm(&b_shader, nir_load_var(&b_shader, vars->stack_ptr), 16),
                              1);
                load_sbt_entry(&b_shader, vars, intr->src[0].ssa, SBT_CALLABLE, 0);
 
-               nir_store_var(
-                  &b_shader, vars->arg,
-                  nir_isub(&b_shader, intr->src[1].ssa, nir_imm_int(&b_shader, size + 16)), 1);
+               nir_store_var(&b_shader, vars->arg,
+                             nir_iadd_imm(&b_shader, intr->src[1].ssa, -size - 16), 1);
 
                vars->stack_sizes[vars->group_idx].recursive_size =
                   MAX2(vars->stack_sizes[vars->group_idx].recursive_size, size + 16);
@@ -402,22 +396,19 @@ lower_rt_instructions(nir_shader *shader, struct rt_variables *vars, unsigned ca
                uint32_t ret = call_idx_base + nir_intrinsic_call_idx(intr) + 1;
                b_shader.cursor = nir_instr_remove(instr);
 
-               nir_store_var(&b_shader, vars->stack_ptr,
-                             nir_iadd(&b_shader, nir_load_var(&b_shader, vars->stack_ptr),
-                                      nir_imm_int(&b_shader, size)),
-                             1);
+               nir_store_var(
+                  &b_shader, vars->stack_ptr,
+                  nir_iadd_imm(&b_shader, nir_load_var(&b_shader, vars->stack_ptr), size), 1);
                nir_store_scratch(&b_shader, nir_imm_int(&b_shader, ret),
                                  nir_load_var(&b_shader, vars->stack_ptr), .align_mul = 16);
 
                nir_store_var(&b_shader, vars->stack_ptr,
-                             nir_iadd(&b_shader, nir_load_var(&b_shader, vars->stack_ptr),
-                                      nir_imm_int(&b_shader, 16)),
+                             nir_iadd_imm(&b_shader, nir_load_var(&b_shader, vars->stack_ptr), 16),
                              1);
 
                nir_store_var(&b_shader, vars->idx, nir_imm_int(&b_shader, 1), 1);
-               nir_store_var(
-                  &b_shader, vars->arg,
-                  nir_isub(&b_shader, intr->src[10].ssa, nir_imm_int(&b_shader, size + 16)), 1);
+               nir_store_var(&b_shader, vars->arg,
+                             nir_iadd_imm(&b_shader, intr->src[10].ssa, -size - 16), 1);
 
                vars->stack_sizes[vars->group_idx].recursive_size =
                   MAX2(vars->stack_sizes[vars->group_idx].recursive_size, size + 16);
@@ -426,17 +417,13 @@ lower_rt_instructions(nir_shader *shader, struct rt_variables *vars, unsigned ca
                nir_store_var(&b_shader, vars->accel_struct, intr->src[0].ssa, 0x1);
                nir_store_var(&b_shader, vars->flags, intr->src[1].ssa, 0x1);
                nir_store_var(&b_shader, vars->cull_mask,
-                             nir_iand(&b_shader, intr->src[2].ssa, nir_imm_int(&b_shader, 0xff)),
-                             0x1);
+                             nir_iand_imm(&b_shader, intr->src[2].ssa, 0xff), 0x1);
                nir_store_var(&b_shader, vars->sbt_offset,
-                             nir_iand(&b_shader, intr->src[3].ssa, nir_imm_int(&b_shader, 0xf)),
-                             0x1);
+                             nir_iand_imm(&b_shader, intr->src[3].ssa, 0xf), 0x1);
                nir_store_var(&b_shader, vars->sbt_stride,
-                             nir_iand(&b_shader, intr->src[4].ssa, nir_imm_int(&b_shader, 0xf)),
-                             0x1);
+                             nir_iand_imm(&b_shader, intr->src[4].ssa, 0xf), 0x1);
                nir_store_var(&b_shader, vars->miss_index,
-                             nir_iand(&b_shader, intr->src[5].ssa, nir_imm_int(&b_shader, 0xffff)),
-                             0x1);
+                             nir_iand_imm(&b_shader, intr->src[5].ssa, 0xffff), 0x1);
                nir_store_var(&b_shader, vars->origin, intr->src[6].ssa, 0x7);
                nir_store_var(&b_shader, vars->tmin, intr->src[7].ssa, 0x1);
                nir_store_var(&b_shader, vars->direction, intr->src[8].ssa, 0x7);
@@ -447,10 +434,9 @@ lower_rt_instructions(nir_shader *shader, struct rt_variables *vars, unsigned ca
                uint32_t size = align(nir_intrinsic_stack_size(intr), 16) + RADV_MAX_HIT_ATTRIB_SIZE;
                b_shader.cursor = nir_instr_remove(instr);
 
-               nir_store_var(&b_shader, vars->stack_ptr,
-                             nir_iadd(&b_shader, nir_load_var(&b_shader, vars->stack_ptr),
-                                      nir_imm_int(&b_shader, -size)),
-                             1);
+               nir_store_var(
+                  &b_shader, vars->stack_ptr,
+                  nir_iadd_imm(&b_shader, nir_load_var(&b_shader, vars->stack_ptr), -size), 1);
                break;
             }
             case nir_intrinsic_rt_return_amd: {
@@ -522,7 +508,7 @@ lower_rt_instructions(nir_shader *shader, struct rt_variables *vars, unsigned ca
             case nir_intrinsic_load_ray_instance_custom_index: {
                b_shader.cursor = nir_instr_remove(instr);
                nir_ssa_def *ret = nir_load_var(&b_shader, vars->custom_instance_and_mask);
-               ret = nir_iand(&b_shader, ret, nir_imm_int(&b_shader, 0xFFFFFF));
+               ret = nir_iand_imm(&b_shader, ret, 0xFFFFFF);
                nir_ssa_def_rewrite_uses(&intr->dest.ssa, ret);
                break;
             }
@@ -535,7 +521,7 @@ lower_rt_instructions(nir_shader *shader, struct rt_variables *vars, unsigned ca
             case nir_intrinsic_load_ray_geometry_index: {
                b_shader.cursor = nir_instr_remove(instr);
                nir_ssa_def *ret = nir_load_var(&b_shader, vars->geometry_id_and_flags);
-               ret = nir_iand(&b_shader, ret, nir_imm_int(&b_shader, 0xFFFFFFF));
+               ret = nir_iand_imm(&b_shader, ret, 0xFFFFFFF);
                nir_ssa_def_rewrite_uses(&intr->dest.ssa, ret);
                break;
             }
@@ -589,9 +575,8 @@ lower_rt_instructions(nir_shader *shader, struct rt_variables *vars, unsigned ca
 
                   val = nir_vec(&b_shader, vals, 3);
                } else {
-                  val = nir_build_load_global(&b_shader, 3, 32,
-                                              nir_iadd(&b_shader, instance_node_addr,
-                                                       nir_imm_int64(&b_shader, 92 + c * 12)));
+                  val = nir_build_load_global(
+                     &b_shader, 3, 32, nir_iadd_imm(&b_shader, instance_node_addr, 92 + c * 12));
                }
                b_shader.cursor = nir_instr_remove(instr);
                nir_ssa_def_rewrite_uses(&intr->dest.ssa, val);
@@ -600,18 +585,15 @@ lower_rt_instructions(nir_shader *shader, struct rt_variables *vars, unsigned ca
             case nir_intrinsic_load_ray_object_origin: {
                nir_ssa_def *instance_node_addr = nir_load_var(&b_shader, vars->instance_addr);
                nir_ssa_def *wto_matrix[] = {
-                  nir_build_load_global(
-                     &b_shader, 4, 32,
-                     nir_iadd(&b_shader, instance_node_addr, nir_imm_int64(&b_shader, 16)),
-                     .align_mul = 64, .align_offset = 16),
-                  nir_build_load_global(
-                     &b_shader, 4, 32,
-                     nir_iadd(&b_shader, instance_node_addr, nir_imm_int64(&b_shader, 32)),
-                     .align_mul = 64, .align_offset = 32),
-                  nir_build_load_global(
-                     &b_shader, 4, 32,
-                     nir_iadd(&b_shader, instance_node_addr, nir_imm_int64(&b_shader, 48)),
-                     .align_mul = 64, .align_offset = 48)};
+                  nir_build_load_global(&b_shader, 4, 32,
+                                        nir_iadd_imm(&b_shader, instance_node_addr, 16),
+                                        .align_mul = 64, .align_offset = 16),
+                  nir_build_load_global(&b_shader, 4, 32,
+                                        nir_iadd_imm(&b_shader, instance_node_addr, 32),
+                                        .align_mul = 64, .align_offset = 32),
+                  nir_build_load_global(&b_shader, 4, 32,
+                                        nir_iadd_imm(&b_shader, instance_node_addr, 48),
+                                        .align_mul = 64, .align_offset = 48)};
                nir_ssa_def *val = nir_build_vec3_mat_mult_pre(
                   &b_shader, nir_load_var(&b_shader, vars->origin), wto_matrix);
                b_shader.cursor = nir_instr_remove(instr);
@@ -718,7 +700,7 @@ insert_rt_case(nir_builder *b, nir_shader *shader, const struct rt_variables *va
          MAX2(src_vars.stack_sizes[src_vars.group_idx].recursive_size, shader->scratch_size);
    }
 
-   nir_push_if(b, nir_ieq(b, idx, nir_imm_int(b, call_idx)));
+   nir_push_if(b, nir_ieq_imm(b, idx, call_idx));
    nir_store_var(b, vars->main_loop_case_visited, nir_imm_bool(b, true), 1);
    nir_inline_function_impl(b, nir_shader_get_entrypoint(shader), NULL, var_remap);
    nir_pop_if(b, NULL);
@@ -1070,7 +1052,7 @@ visit_any_hit_shaders(struct radv_device *device,
 {
    nir_ssa_def *sbt_idx = nir_load_var(b, vars->idx);
 
-   nir_push_if(b, nir_ine(b, sbt_idx, nir_imm_int(b, 0)));
+   nir_push_if(b, nir_ine_imm(b, sbt_idx, 0));
    for (unsigned i = 0; i < pCreateInfo->groupCount; ++i) {
       const VkRayTracingShaderGroupCreateInfoKHR *group_info = &pCreateInfo->pGroups[i];
       uint32_t shader_id = VK_SHADER_UNUSED_KHR;
@@ -1104,34 +1086,30 @@ insert_traversal_triangle_case(struct radv_device *device,
    nir_ssa_def *div = nir_vector_extract(b, result, nir_imm_int(b, 1));
    dist = nir_fdiv(b, dist, div);
    nir_ssa_def *frontface = nir_flt(b, nir_imm_float(b, 0), div);
-   nir_ssa_def *switch_ccw = nir_ine(
-      b,
-      nir_iand(
-         b, nir_load_var(b, trav_vars->sbt_offset_and_flags),
-         nir_imm_int(b, VK_GEOMETRY_INSTANCE_TRIANGLE_FRONT_COUNTERCLOCKWISE_BIT_KHR << 24)),
-      nir_imm_int(b, 0));
+   nir_ssa_def *switch_ccw =
+      nir_ine_imm(b,
+                  nir_iand_imm(b, nir_load_var(b, trav_vars->sbt_offset_and_flags),
+                               VK_GEOMETRY_INSTANCE_TRIANGLE_FRONT_COUNTERCLOCKWISE_BIT_KHR << 24),
+                  0);
    frontface = nir_ixor(b, frontface, switch_ccw);
 
-   nir_ssa_def *not_cull = nir_ieq(
-      b, nir_iand(b, nir_load_var(b, vars->flags), nir_imm_int(b, SpvRayFlagsSkipTrianglesKHRMask)),
-      nir_imm_int(b, 0));
-   nir_ssa_def *not_facing_cull = nir_ieq(
+   nir_ssa_def *not_cull = nir_ieq_imm(
+      b, nir_iand_imm(b, nir_load_var(b, vars->flags), SpvRayFlagsSkipTrianglesKHRMask), 0);
+   nir_ssa_def *not_facing_cull = nir_ieq_imm(
       b,
       nir_iand(b, nir_load_var(b, vars->flags),
                nir_bcsel(b, frontface, nir_imm_int(b, SpvRayFlagsCullFrontFacingTrianglesKHRMask),
                          nir_imm_int(b, SpvRayFlagsCullBackFacingTrianglesKHRMask))),
-      nir_imm_int(b, 0));
+      0);
 
    not_cull = nir_iand(
       b, not_cull,
       nir_ior(
          b, not_facing_cull,
-         nir_ine(
-            b,
-            nir_iand(
-               b, nir_load_var(b, trav_vars->sbt_offset_and_flags),
-               nir_imm_int(b, VK_GEOMETRY_INSTANCE_TRIANGLE_FACING_CULL_DISABLE_BIT_KHR << 24)),
-            nir_imm_int(b, 0))));
+         nir_ine_imm(b,
+                     nir_iand_imm(b, nir_load_var(b, trav_vars->sbt_offset_and_flags),
+                                  VK_GEOMETRY_INSTANCE_TRIANGLE_FACING_CULL_DISABLE_BIT_KHR << 24),
+                     0)));
 
    nir_push_if(b, nir_iand(b,
                            nir_iand(b, nir_flt(b, dist, nir_load_var(b, vars->tmax)),
@@ -1139,38 +1117,36 @@ insert_traversal_triangle_case(struct radv_device *device,
                            not_cull));
    {
 
-      nir_ssa_def *triangle_info = nir_build_load_global(
-         b, 2, 32,
-         nir_iadd(b, build_node_to_addr(device, b, bvh_node),
-                  nir_imm_int64(b, offsetof(struct radv_bvh_triangle_node, triangle_id))));
+      nir_ssa_def *triangle_info =
+         nir_build_load_global(b, 2, 32,
+                               nir_iadd_imm(b, build_node_to_addr(device, b, bvh_node),
+                                            offsetof(struct radv_bvh_triangle_node, triangle_id)));
       nir_ssa_def *primitive_id = nir_channel(b, triangle_info, 0);
       nir_ssa_def *geometry_id_and_flags = nir_channel(b, triangle_info, 1);
-      nir_ssa_def *geometry_id = nir_iand(b, geometry_id_and_flags, nir_imm_int(b, 0xfffffff));
+      nir_ssa_def *geometry_id = nir_iand_imm(b, geometry_id_and_flags, 0xfffffff);
       nir_ssa_def *is_opaque = hit_is_opaque(b, nir_load_var(b, trav_vars->sbt_offset_and_flags),
                                              nir_load_var(b, vars->flags), geometry_id_and_flags);
 
       not_cull =
-         nir_ieq(b,
-                 nir_iand(b, nir_load_var(b, vars->flags),
-                          nir_bcsel(b, is_opaque, nir_imm_int(b, SpvRayFlagsCullOpaqueKHRMask),
-                                    nir_imm_int(b, SpvRayFlagsCullNoOpaqueKHRMask))),
-                 nir_imm_int(b, 0));
+         nir_ieq_imm(b,
+                     nir_iand(b, nir_load_var(b, vars->flags),
+                              nir_bcsel(b, is_opaque, nir_imm_int(b, SpvRayFlagsCullOpaqueKHRMask),
+                                        nir_imm_int(b, SpvRayFlagsCullNoOpaqueKHRMask))),
+                     0);
       nir_push_if(b, not_cull);
       {
-         nir_ssa_def *sbt_idx =
-            nir_iadd(b,
-                     nir_iadd(b, nir_load_var(b, vars->sbt_offset),
-                              nir_iand(b, nir_load_var(b, trav_vars->sbt_offset_and_flags),
-                                       nir_imm_int(b, 0xffffff))),
-                     nir_imul(b, nir_load_var(b, vars->sbt_stride), geometry_id));
+         nir_ssa_def *sbt_idx = nir_iadd(
+            b,
+            nir_iadd(b, nir_load_var(b, vars->sbt_offset),
+                     nir_iand_imm(b, nir_load_var(b, trav_vars->sbt_offset_and_flags), 0xffffff)),
+            nir_imul(b, nir_load_var(b, vars->sbt_stride), geometry_id));
          nir_ssa_def *divs[2] = {div, div};
          nir_ssa_def *ij = nir_fdiv(b, nir_channels(b, result, 0xc), nir_vec(b, divs, 2));
          nir_ssa_def *hit_kind =
             nir_bcsel(b, frontface, nir_imm_int(b, 0xFE), nir_imm_int(b, 0xFF));
 
          nir_store_scratch(
-            b, ij,
-            nir_iadd(b, nir_load_var(b, vars->stack_ptr), nir_imm_int(b, RADV_HIT_ATTRIB_OFFSET)),
+            b, ij, nir_iadd_imm(b, nir_load_var(b, vars->stack_ptr), RADV_HIT_ATTRIB_OFFSET),
             .align_mul = 16);
 
          nir_store_var(b, vars->ahit_status, nir_imm_int(b, 0), 1);
@@ -1193,7 +1169,7 @@ insert_traversal_triangle_case(struct radv_device *device,
 
             visit_any_hit_shaders(device, pCreateInfo, b, &inner_vars);
 
-            nir_push_if(b, nir_ieq(b, nir_load_var(b, vars->ahit_status), nir_imm_int(b, 1)));
+            nir_push_if(b, nir_ieq_imm(b, nir_load_var(b, vars->ahit_status), 1));
             {
                nir_jump(b, nir_jump_continue);
             }
@@ -1214,20 +1190,17 @@ insert_traversal_triangle_case(struct radv_device *device,
 
          nir_store_var(b, trav_vars->should_return,
                        nir_ior(b,
-                               nir_ine(b,
-                                       nir_iand(b, nir_load_var(b, vars->flags),
-                                                nir_imm_int(b, SpvRayFlagsSkipClosestHitShaderKHRMask)),
-                                       nir_imm_int(b, 0)),
-                               nir_ieq(b, nir_load_var(b, vars->idx), nir_imm_int(b, 0))),
+                               nir_ine_imm(b,
+                                           nir_iand_imm(b, nir_load_var(b, vars->flags),
+                                                        SpvRayFlagsSkipClosestHitShaderKHRMask),
+                                           0),
+                               nir_ieq_imm(b, nir_load_var(b, vars->idx), 0)),
                        1);
 
-         nir_ssa_def *terminate_on_first_hit =
-            nir_ine(b,
-                    nir_iand(b, nir_load_var(b, vars->flags),
-                             nir_imm_int(b, SpvRayFlagsTerminateOnFirstHitKHRMask)),
-                    nir_imm_int(b, 0));
-         nir_ssa_def *ray_terminated =
-            nir_ieq(b, nir_load_var(b, vars->ahit_status), nir_imm_int(b, 2));
+         nir_ssa_def *terminate_on_first_hit = nir_ine_imm(
+            b, nir_iand_imm(b, nir_load_var(b, vars->flags), SpvRayFlagsTerminateOnFirstHitKHRMask),
+            0);
+         nir_ssa_def *ray_terminated = nir_ieq_imm(b, nir_load_var(b, vars->ahit_status), 2);
          nir_push_if(b, nir_ior(b, terminate_on_first_hit, ray_terminated));
          {
             nir_jump(b, nir_jump_break);
@@ -1246,31 +1219,29 @@ insert_traversal_aabb_case(struct radv_device *device,
                            const struct rt_traversal_vars *trav_vars, nir_ssa_def *bvh_node)
 {
    nir_ssa_def *node_addr = build_node_to_addr(device, b, bvh_node);
-   nir_ssa_def *triangle_info =
-      nir_build_load_global(b, 2, 32, nir_iadd(b, node_addr, nir_imm_int64(b, 24)));
+   nir_ssa_def *triangle_info = nir_build_load_global(b, 2, 32, nir_iadd_imm(b, node_addr, 24));
    nir_ssa_def *primitive_id = nir_channel(b, triangle_info, 0);
    nir_ssa_def *geometry_id_and_flags = nir_channel(b, triangle_info, 1);
-   nir_ssa_def *geometry_id = nir_iand(b, geometry_id_and_flags, nir_imm_int(b, 0xfffffff));
+   nir_ssa_def *geometry_id = nir_iand_imm(b, geometry_id_and_flags, 0xfffffff);
    nir_ssa_def *is_opaque = hit_is_opaque(b, nir_load_var(b, trav_vars->sbt_offset_and_flags),
                                           nir_load_var(b, vars->flags), geometry_id_and_flags);
 
-   nir_ssa_def *not_skip_aabb = nir_ieq(
-      b, nir_iand(b, nir_load_var(b, vars->flags), nir_imm_int(b, SpvRayFlagsSkipAABBsKHRMask)),
-      nir_imm_int(b, 0));
-   nir_ssa_def *not_cull =
-      nir_iand(b, not_skip_aabb, nir_ieq(b,
-              nir_iand(b, nir_load_var(b, vars->flags),
-                       nir_bcsel(b, is_opaque, nir_imm_int(b, SpvRayFlagsCullOpaqueKHRMask),
-                                 nir_imm_int(b, SpvRayFlagsCullNoOpaqueKHRMask))),
-              nir_imm_int(b, 0)));
+   nir_ssa_def *not_skip_aabb =
+      nir_ieq_imm(b, nir_iand_imm(b, nir_load_var(b, vars->flags), SpvRayFlagsSkipAABBsKHRMask), 0);
+   nir_ssa_def *not_cull = nir_iand(
+      b, not_skip_aabb,
+      nir_ieq_imm(b,
+                  nir_iand(b, nir_load_var(b, vars->flags),
+                           nir_bcsel(b, is_opaque, nir_imm_int(b, SpvRayFlagsCullOpaqueKHRMask),
+                                     nir_imm_int(b, SpvRayFlagsCullNoOpaqueKHRMask))),
+                  0));
    nir_push_if(b, not_cull);
    {
-      nir_ssa_def *sbt_idx =
-         nir_iadd(b,
-                  nir_iadd(b, nir_load_var(b, vars->sbt_offset),
-                           nir_iand(b, nir_load_var(b, trav_vars->sbt_offset_and_flags),
-                                    nir_imm_int(b, 0xffffff))),
-                  nir_imul(b, nir_load_var(b, vars->sbt_stride), geometry_id));
+      nir_ssa_def *sbt_idx = nir_iadd(
+         b,
+         nir_iadd(b, nir_load_var(b, vars->sbt_offset),
+                  nir_iand_imm(b, nir_load_var(b, trav_vars->sbt_offset_and_flags), 0xffffff)),
+         nir_imul(b, nir_load_var(b, vars->sbt_stride), geometry_id));
 
       struct rt_variables inner_vars = create_inner_vars(b, vars);
 
@@ -1291,7 +1262,7 @@ insert_traversal_aabb_case(struct radv_device *device,
 
       nir_store_var(b, vars->ahit_status, nir_imm_int(b, 1), 1);
 
-      nir_push_if(b, nir_ine(b, nir_load_var(b, inner_vars.idx), nir_imm_int(b, 0)));
+      nir_push_if(b, nir_ine_imm(b, nir_load_var(b, inner_vars.idx), 0));
       for (unsigned i = 0; i < pCreateInfo->groupCount; ++i) {
          const VkRayTracingShaderGroupCreateInfoKHR *group_info = &pCreateInfo->pGroups[i];
          uint32_t shader_id = VK_SHADER_UNUSED_KHR;
@@ -1329,10 +1300,8 @@ insert_traversal_aabb_case(struct radv_device *device,
          nir_ssa_def *vec3_inf =
             nir_channels(b, nir_imm_vec4(b, INFINITY, INFINITY, INFINITY, 0), 0x7);
 
-         nir_ssa_def *bvh_lo =
-            nir_build_load_global(b, 3, 32, nir_iadd(b, node_addr, nir_imm_int64(b, 0)));
-         nir_ssa_def *bvh_hi =
-            nir_build_load_global(b, 3, 32, nir_iadd(b, node_addr, nir_imm_int64(b, 12)));
+         nir_ssa_def *bvh_lo = nir_build_load_global(b, 3, 32, nir_iadd_imm(b, node_addr, 0));
+         nir_ssa_def *bvh_hi = nir_build_load_global(b, 3, 32, nir_iadd_imm(b, node_addr, 12));
 
          bvh_lo = nir_fsub(b, bvh_lo, nir_load_var(b, trav_vars->origin));
          bvh_hi = nir_fsub(b, bvh_hi, nir_load_var(b, trav_vars->origin));
@@ -1360,7 +1329,7 @@ insert_traversal_aabb_case(struct radv_device *device,
       }
       nir_pop_if(b, NULL);
 
-      nir_push_if(b, nir_ine(b, nir_load_var(b, vars->ahit_status), nir_imm_int(b, 1)));
+      nir_push_if(b, nir_ine_imm(b, nir_load_var(b, vars->ahit_status), 1));
       {
          nir_store_var(b, vars->primitive_id, primitive_id, 1);
          nir_store_var(b, vars->geometry_id_and_flags, geometry_id_and_flags, 1);
@@ -1374,20 +1343,17 @@ insert_traversal_aabb_case(struct radv_device *device,
 
          nir_store_var(b, trav_vars->should_return,
                        nir_ior(b,
-                               nir_ine(b,
-                                       nir_iand(b, nir_load_var(b, vars->flags),
-                                                nir_imm_int(b, SpvRayFlagsSkipClosestHitShaderKHRMask)),
-                                       nir_imm_int(b, 0)),
-                               nir_ieq(b, nir_load_var(b, vars->idx), nir_imm_int(b, 0))),
+                               nir_ine_imm(b,
+                                           nir_iand_imm(b, nir_load_var(b, vars->flags),
+                                                        SpvRayFlagsSkipClosestHitShaderKHRMask),
+                                           0),
+                               nir_ieq_imm(b, nir_load_var(b, vars->idx), 0)),
                        1);
 
-         nir_ssa_def *terminate_on_first_hit =
-            nir_ine(b,
-                    nir_iand(b, nir_load_var(b, vars->flags),
-                             nir_imm_int(b, SpvRayFlagsTerminateOnFirstHitKHRMask)),
-                    nir_imm_int(b, 0));
-         nir_ssa_def *ray_terminated =
-            nir_ieq(b, nir_load_var(b, vars->ahit_status), nir_imm_int(b, 2));
+         nir_ssa_def *terminate_on_first_hit = nir_ine_imm(
+            b, nir_iand_imm(b, nir_load_var(b, vars->flags), SpvRayFlagsTerminateOnFirstHitKHRMask),
+            0);
+         nir_ssa_def *ray_terminated = nir_ieq_imm(b, nir_load_var(b, vars->ahit_status), 2);
          nir_push_if(b, nir_ior(b, terminate_on_first_hit, ray_terminated));
          {
             nir_jump(b, nir_jump_break);
@@ -1409,8 +1375,8 @@ insert_traversal(struct radv_device *device, const VkRayTracingPipelineCreateInf
    unsigned stack_entry_stride = stack_entry_size * lanes;
    nir_ssa_def *stack_entry_stride_def = nir_imm_int(b, stack_entry_stride);
    nir_ssa_def *stack_base =
-      nir_iadd(b, nir_imm_int(b, b->shader->info.shared_size),
-               nir_imul(b, nir_load_local_invocation_index(b), nir_imm_int(b, stack_entry_size)));
+      nir_iadd_imm(b, nir_imul_imm(b, nir_load_local_invocation_index(b), stack_entry_size),
+                   b->shader->info.shared_size);
 
    b->shader->info.shared_size += stack_entry_stride * MAX_STACK_ENTRY_COUNT;
    assert(b->shader->info.shared_size <= 32768);
@@ -1425,7 +1391,7 @@ insert_traversal(struct radv_device *device, const VkRayTracingPipelineCreateInf
 
    nir_store_var(b, trav_vars.should_return, nir_imm_bool(b, false), 1);
 
-   nir_push_if(b, nir_ine(b, accel_struct, nir_imm_int64(b, 0)));
+   nir_push_if(b, nir_ine_imm(b, accel_struct, 0));
    {
       nir_store_var(b, trav_vars.bvh_base, build_addr_to_node(b, accel_struct), 1);
 
@@ -1469,7 +1435,7 @@ insert_traversal(struct radv_device *device, const VkRayTracingPipelineCreateInf
 
       nir_ssa_def *bvh_node = nir_load_shared(b, 1, 32, nir_load_var(b, trav_vars.stack), .base = 0,
                                               .align_mul = stack_entry_size);
-      nir_ssa_def *bvh_node_type = nir_iand(b, bvh_node, nir_imm_int(b, 7));
+      nir_ssa_def *bvh_node_type = nir_iand_imm(b, bvh_node, 7);
 
       bvh_node = nir_iadd(b, nir_load_var(b, trav_vars.bvh_base), nir_u2u(b, bvh_node, 64));
       nir_ssa_def *intrinsic_result = NULL;
@@ -1480,14 +1446,12 @@ insert_traversal(struct radv_device *device, const VkRayTracingPipelineCreateInf
             nir_load_var(b, trav_vars.inv_dir));
       }
 
-      nir_push_if(b, nir_ine(b, nir_iand(b, bvh_node_type, nir_imm_int(b, 4)), nir_imm_int(b, 0)));
+      nir_push_if(b, nir_ine_imm(b, nir_iand_imm(b, bvh_node_type, 4), 0));
       {
-         nir_push_if(b,
-                     nir_ine(b, nir_iand(b, bvh_node_type, nir_imm_int(b, 2)), nir_imm_int(b, 0)));
+         nir_push_if(b, nir_ine_imm(b, nir_iand_imm(b, bvh_node_type, 2), 0));
          {
             /* custom */
-            nir_push_if(
-               b, nir_ine(b, nir_iand(b, bvh_node_type, nir_imm_int(b, 1)), nir_imm_int(b, 0)));
+            nir_push_if(b, nir_ine_imm(b, nir_iand_imm(b, bvh_node_type, 1), 0));
             if (!(pCreateInfo->flags & VK_PIPELINE_CREATE_RAY_TRACING_SKIP_AABBS_BIT_KHR)) {
                insert_traversal_aabb_case(device, pCreateInfo, b, vars, &trav_vars, bvh_node);
             }
@@ -1498,23 +1462,20 @@ insert_traversal(struct radv_device *device, const VkRayTracingPipelineCreateInf
                nir_ssa_def *instance_data =
                   nir_build_load_global(b, 4, 32, instance_node_addr, .align_mul = 64);
                nir_ssa_def *wto_matrix[] = {
-                  nir_build_load_global(b, 4, 32,
-                                        nir_iadd(b, instance_node_addr, nir_imm_int64(b, 16)),
+                  nir_build_load_global(b, 4, 32, nir_iadd_imm(b, instance_node_addr, 16),
                                         .align_mul = 64, .align_offset = 16),
-                  nir_build_load_global(b, 4, 32,
-                                        nir_iadd(b, instance_node_addr, nir_imm_int64(b, 32)),
+                  nir_build_load_global(b, 4, 32, nir_iadd_imm(b, instance_node_addr, 32),
                                         .align_mul = 64, .align_offset = 32),
-                  nir_build_load_global(b, 4, 32,
-                                        nir_iadd(b, instance_node_addr, nir_imm_int64(b, 48)),
+                  nir_build_load_global(b, 4, 32, nir_iadd_imm(b, instance_node_addr, 48),
                                         .align_mul = 64, .align_offset = 48)};
-               nir_ssa_def *instance_id = nir_build_load_global(
-                  b, 1, 32, nir_iadd(b, instance_node_addr, nir_imm_int64(b, 88)));
+               nir_ssa_def *instance_id =
+                  nir_build_load_global(b, 1, 32, nir_iadd_imm(b, instance_node_addr, 88));
                nir_ssa_def *instance_and_mask = nir_channel(b, instance_data, 2);
-               nir_ssa_def *instance_mask = nir_ushr(b, instance_and_mask, nir_imm_int(b, 24));
+               nir_ssa_def *instance_mask = nir_ushr_imm(b, instance_and_mask, 24);
 
-               nir_push_if(b,
-                           nir_ieq(b, nir_iand(b, instance_mask, nir_load_var(b, vars->cull_mask)),
-                                   nir_imm_int(b, 0)));
+               nir_push_if(
+                  b,
+                  nir_ieq_imm(b, nir_iand(b, instance_mask, nir_load_var(b, vars->cull_mask)), 0));
                nir_jump(b, nir_jump_continue);
                nir_pop_if(b, NULL);
 
@@ -1523,9 +1484,9 @@ insert_traversal(struct radv_device *device, const VkRayTracingPipelineCreateInf
                              build_addr_to_node(
                                 b, nir_pack_64_2x32(b, nir_channels(b, instance_data, 0x3))),
                              1);
-               nir_store_shared(
-                  b, nir_iand(b, nir_channel(b, instance_data, 0), nir_imm_int(b, 63)),
-                  nir_load_var(b, trav_vars.stack), .base = 0, .align_mul = stack_entry_size);
+               nir_store_shared(b, nir_iand_imm(b, nir_channel(b, instance_data, 0), 63),
+                                nir_load_var(b, trav_vars.stack), .base = 0,
+                                .align_mul = stack_entry_size);
                nir_store_var(b, trav_vars.stack,
                              nir_iadd(b, nir_load_var(b, trav_vars.stack), stack_entry_stride_def),
                              1);
@@ -1561,7 +1522,7 @@ insert_traversal(struct radv_device *device, const VkRayTracingPipelineCreateInf
 
             for (unsigned i = 4; i-- > 0; ) {
                nir_ssa_def *new_node = nir_vector_extract(b, result, nir_imm_int(b, i));
-               nir_push_if(b, nir_ine(b, new_node, nir_imm_int(b, 0xffffffff)));
+               nir_push_if(b, nir_ine_imm(b, new_node, 0xffffffff));
                {
                   nir_store_shared(b, new_node, nir_load_var(b, trav_vars.stack), .base = 0,
                                    .align_mul = stack_entry_size);
@@ -1603,7 +1564,7 @@ insert_traversal(struct radv_device *device, const VkRayTracingPipelineCreateInf
       /* Only load the miss shader if we actually miss, which we determining by not having set
        * a closest hit shader. It is valid to not specify an SBT pointer for miss shaders if none
        * of the rays miss. */
-      nir_push_if(b, nir_ieq(b, nir_load_var(b, vars->idx), nir_imm_int(b, 0)));
+      nir_push_if(b, nir_ieq_imm(b, nir_load_var(b, vars->idx), 0));
       {
          load_sbt_entry(b, vars, nir_load_var(b, vars->miss_index), SBT_MISS, 0);
       }
@@ -1704,7 +1665,7 @@ create_rt_shader(struct radv_device *device, const VkRayTracingPipelineCreateInf
 
    nir_loop *loop = nir_push_loop(&b);
 
-   nir_push_if(&b, nir_ior(&b, nir_ieq(&b, nir_load_var(&b, vars.idx), nir_imm_int(&b, 0)),
+   nir_push_if(&b, nir_ior(&b, nir_ieq_imm(&b, nir_load_var(&b, vars.idx), 0),
                            nir_ine(&b, nir_load_var(&b, vars.main_loop_case_visited),
                                    nir_imm_bool(&b, true))));
    nir_jump(&b, nir_jump_break);
@@ -1712,7 +1673,7 @@ create_rt_shader(struct radv_device *device, const VkRayTracingPipelineCreateInf
 
    nir_store_var(&b, vars.main_loop_case_visited, nir_imm_bool(&b, false), 1);
 
-   nir_push_if(&b, nir_ieq(&b, nir_load_var(&b, vars.idx), nir_imm_int(&b, 1)));
+   nir_push_if(&b, nir_ieq_imm(&b, nir_load_var(&b, vars.idx), 1));
    nir_store_var(&b, vars.main_loop_case_visited, nir_imm_bool(&b, true), 1);
    insert_traversal(device, pCreateInfo, &b, &vars);
    nir_pop_if(&b, NULL);
diff --git a/src/amd/vulkan/radv_query.c b/src/amd/vulkan/radv_query.c
index 22ba3a70810..f4e11e903f8 100644
--- a/src/amd/vulkan/radv_query.c
+++ b/src/amd/vulkan/radv_query.c
@@ -52,7 +52,7 @@ radv_get_pipeline_statistics_index(const VkQueryPipelineStatisticFlagBits flag)
 static nir_ssa_def *
 nir_test_flag(nir_builder *b, nir_ssa_def *flags, uint32_t flag)
 {
-   return nir_i2b(b, nir_iand(b, flags, nir_imm_int(b, flag)));
+   return nir_i2b(b, nir_iand_imm(b, flags, flag));
 }
 
 static void
@@ -149,12 +149,12 @@ build_occlusion_query_shader(struct radv_device *device)
    nir_ssa_def *current_outer_count = nir_load_var(&b, outer_counter);
    radv_break_on_count(&b, outer_counter, nir_imm_int(&b, db_count));
 
-   nir_ssa_def *enabled_cond = nir_iand(&b, nir_imm_int(&b, enabled_rb_mask),
-                                        nir_ishl(&b, nir_imm_int(&b, 1), current_outer_count));
+   nir_ssa_def *enabled_cond =
+      nir_iand_imm(&b, nir_ishl(&b, nir_imm_int(&b, 1), current_outer_count), enabled_rb_mask);
 
    nir_push_if(&b, nir_i2b(&b, enabled_cond));
 
-   nir_ssa_def *load_offset = nir_imul(&b, current_outer_count, nir_imm_int(&b, 16));
+   nir_ssa_def *load_offset = nir_imul_imm(&b, current_outer_count, 16);
    load_offset = nir_iadd(&b, input_base, load_offset);
 
    nir_ssa_def *load = nir_load_ssbo(&b, 2, 64, src_buf, load_offset, .align_mul = 16);
@@ -271,13 +271,13 @@ build_pipeline_statistics_query_shader(struct radv_device *device)
    nir_ssa_def *output_stride = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 4), .range = 16);
    nir_ssa_def *output_base = nir_imul(&b, output_stride, global_id);
 
-   avail_offset = nir_iadd(&b, avail_offset, nir_imul(&b, global_id, nir_imm_int(&b, 4)));
+   avail_offset = nir_iadd(&b, avail_offset, nir_imul_imm(&b, global_id, 4));
 
    nir_ssa_def *available32 = nir_load_ssbo(&b, 1, 32, src_buf, avail_offset);
 
    nir_ssa_def *result_is_64bit = nir_test_flag(&b, flags, VK_QUERY_RESULT_64_BIT);
    nir_ssa_def *elem_size = nir_bcsel(&b, result_is_64bit, nir_imm_int(&b, 8), nir_imm_int(&b, 4));
-   nir_ssa_def *elem_count = nir_ushr(&b, stats_mask, nir_imm_int(&b, 16));
+   nir_ssa_def *elem_count = nir_ushr_imm(&b, stats_mask, 16);
 
    radv_store_availability(&b, flags, dst_buf,
                            nir_iadd(&b, output_base, nir_imul(&b, elem_count, elem_size)),
@@ -289,13 +289,11 @@ build_pipeline_statistics_query_shader(struct radv_device *device)
    for (int i = 0; i < ARRAY_SIZE(pipeline_statistics_indices); ++i) {
       nir_push_if(&b, nir_test_flag(&b, stats_mask, 1u << i));
 
-      nir_ssa_def *start_offset =
-         nir_iadd(&b, input_base, nir_imm_int(&b, pipeline_statistics_indices[i] * 8));
+      nir_ssa_def *start_offset = nir_iadd_imm(&b, input_base, pipeline_statistics_indices[i] * 8);
       nir_ssa_def *start = nir_load_ssbo(&b, 1, 64, src_buf, start_offset);
 
       nir_ssa_def *end_offset =
-         nir_iadd(&b, input_base,
-                  nir_imm_int(&b, pipeline_statistics_indices[i] * 8 + pipelinestat_block_size));
+         nir_iadd_imm(&b, input_base, pipeline_statistics_indices[i] * 8 + pipelinestat_block_size);
       nir_ssa_def *end = nir_load_ssbo(&b, 1, 64, src_buf, end_offset);
 
       nir_ssa_def *result = nir_isub(&b, end, start);
@@ -414,15 +412,15 @@ build_tfb_query_shader(struct radv_device *device)
 
    /* Load data from the query pool. */
    nir_ssa_def *load1 = nir_load_ssbo(&b, 4, 32, src_buf, input_base, .align_mul = 32);
-   nir_ssa_def *load2 = nir_load_ssbo(
-      &b, 4, 32, src_buf, nir_iadd(&b, input_base, nir_imm_int(&b, 16)), .align_mul = 16);
+   nir_ssa_def *load2 =
+      nir_load_ssbo(&b, 4, 32, src_buf, nir_iadd_imm(&b, input_base, 16), .align_mul = 16);
 
    /* Check if result is available. */
    nir_ssa_def *avails[2];
    avails[0] = nir_iand(&b, nir_channel(&b, load1, 1), nir_channel(&b, load1, 3));
    avails[1] = nir_iand(&b, nir_channel(&b, load2, 1), nir_channel(&b, load2, 3));
    nir_ssa_def *result_is_available =
-      nir_i2b(&b, nir_iand(&b, nir_iand(&b, avails[0], avails[1]), nir_imm_int(&b, 0x80000000)));
+      nir_i2b(&b, nir_iand_imm(&b, nir_iand(&b, avails[0], avails[1]), 0x80000000));
 
    /* Only compute result if available. */
    nir_push_if(&b, result_is_available);
@@ -541,8 +539,7 @@ build_timestamp_query_shader(struct radv_device *device)
       nir_pack_64_2x32(&b, nir_vec2(&b, nir_channel(&b, load, 0), nir_channel(&b, load, 1)));
 
    /* Check if result is available. */
-   nir_ssa_def *result_is_available =
-      nir_i2b(&b, nir_ine(&b, timestamp, nir_imm_int64(&b, TIMESTAMP_NOT_READY)));
+   nir_ssa_def *result_is_available = nir_i2b(&b, nir_ine_imm(&b, timestamp, TIMESTAMP_NOT_READY));
 
    /* Only store result if available. */
    nir_push_if(&b, result_is_available);
diff --git a/src/amd/vulkan/radv_rt_common.c b/src/amd/vulkan/radv_rt_common.c
index d30dbf21a88..e12bb300d66 100644
--- a/src/amd/vulkan/radv_rt_common.c
+++ b/src/amd/vulkan/radv_rt_common.c
@@ -101,13 +101,13 @@ intersect_ray_amd_software_box(struct radv_device *device, nir_builder *b, nir_s
 
       /* node->children[i] -> uint */
       nir_ssa_def *child_index =
-         nir_build_load_global(b, 1, 32, nir_iadd(b, node_addr, nir_imm_int64(b, child_offset)),
-                               .align_mul = 64, .align_offset = child_offset % 64);
+         nir_build_load_global(b, 1, 32, nir_iadd_imm(b, node_addr, child_offset), .align_mul = 64,
+                               .align_offset = child_offset % 64);
       /* node->coords[i][0], node->coords[i][1] -> vec3 */
       nir_ssa_def *node_coords[2] = {
-         nir_build_load_global(b, 3, 32, nir_iadd(b, node_addr, nir_imm_int64(b, coord_offsets[0])),
+         nir_build_load_global(b, 3, 32, nir_iadd_imm(b, node_addr, coord_offsets[0]),
                                .align_mul = 64, .align_offset = coord_offsets[0] % 64),
-         nir_build_load_global(b, 3, 32, nir_iadd(b, node_addr, nir_imm_int64(b, coord_offsets[1])),
+         nir_build_load_global(b, 3, 32, nir_iadd_imm(b, node_addr, coord_offsets[1]),
                                .align_mul = 64, .align_offset = coord_offsets[1] % 64),
       };
 
@@ -185,12 +185,12 @@ intersect_ray_amd_software_tri(struct radv_device *device, nir_builder *b, nir_s
 
    /* node->coords[0], node->coords[1], node->coords[2] -> vec3 */
    nir_ssa_def *node_coords[3] = {
-      nir_build_load_global(b, 3, 32, nir_iadd(b, node_addr, nir_imm_int64(b, coord_offsets[0])),
-                            .align_mul = 64, .align_offset = coord_offsets[0] % 64),
-      nir_build_load_global(b, 3, 32, nir_iadd(b, node_addr, nir_imm_int64(b, coord_offsets[1])),
-                            .align_mul = 64, .align_offset = coord_offsets[1] % 64),
-      nir_build_load_global(b, 3, 32, nir_iadd(b, node_addr, nir_imm_int64(b, coord_offsets[2])),
-                            .align_mul = 64, .align_offset = coord_offsets[2] % 64),
+      nir_build_load_global(b, 3, 32, nir_iadd_imm(b, node_addr, coord_offsets[0]), .align_mul = 64,
+                            .align_offset = coord_offsets[0] % 64),
+      nir_build_load_global(b, 3, 32, nir_iadd_imm(b, node_addr, coord_offsets[1]), .align_mul = 64,
+                            .align_offset = coord_offsets[1] % 64),
+      nir_build_load_global(b, 3, 32, nir_iadd_imm(b, node_addr, coord_offsets[2]), .align_mul = 64,
+                            .align_offset = coord_offsets[2] % 64),
    };
 
    nir_variable *result = nir_variable_create(b->shader, nir_var_shader_temp, vec4_type, "result");
@@ -212,8 +212,8 @@ intersect_ray_amd_software_tri(struct radv_device *device, nir_builder *b, nir_s
       b, nir_fge(b, abs_dirs[0], abs_dirs[1]),
       nir_bcsel(b, nir_fge(b, abs_dirs[0], abs_dirs[2]), nir_imm_int(b, 0), nir_imm_int(b, 2)),
       nir_bcsel(b, nir_fge(b, abs_dirs[1], abs_dirs[2]), nir_imm_int(b, 1), nir_imm_int(b, 2)));
-   nir_ssa_def *kx = nir_imod(b, nir_iadd(b, kz, nir_imm_int(b, 1)), nir_imm_int(b, 3));
-   nir_ssa_def *ky = nir_imod(b, nir_iadd(b, kx, nir_imm_int(b, 1)), nir_imm_int(b, 3));
+   nir_ssa_def *kx = nir_imod(b, nir_iadd_imm(b, kz, 1), nir_imm_int(b, 3));
+   nir_ssa_def *ky = nir_imod(b, nir_iadd_imm(b, kx, 1), nir_imm_int(b, 3));
    nir_ssa_def *k_indices[3] = {kx, ky, kz};
    nir_ssa_def *k = nir_vec(b, k_indices, 3);
 
@@ -337,19 +337,19 @@ nir_ssa_def *
 build_addr_to_node(nir_builder *b, nir_ssa_def *addr)
 {
    const uint64_t bvh_size = 1ull << 42;
-   nir_ssa_def *node = nir_ushr(b, addr, nir_imm_int(b, 3));
-   return nir_iand(b, node, nir_imm_int64(b, (bvh_size - 1) << 3));
+   nir_ssa_def *node = nir_ushr_imm(b, addr, 3);
+   return nir_iand_imm(b, node, (bvh_size - 1) << 3);
 }
 
 nir_ssa_def *
 build_node_to_addr(struct radv_device *device, nir_builder *b, nir_ssa_def *node)
 {
-   nir_ssa_def *addr = nir_iand(b, node, nir_imm_int64(b, ~7ull));
-   addr = nir_ishl(b, addr, nir_imm_int(b, 3));
+   nir_ssa_def *addr = nir_iand_imm(b, node, ~7ull);
+   addr = nir_ishl_imm(b, addr, 3);
    /* Assumes everything is in the top half of address space, which is true in
     * GFX9+ for now. */
    return device->physical_device->rad_info.chip_class >= GFX9
-             ? nir_ior(b, addr, nir_imm_int64(b, 0xffffull << 48))
+             ? nir_ior_imm(b, addr, 0xffffull << 48)
              : addr;
 }
 
@@ -388,8 +388,7 @@ nir_build_wto_matrix_load(nir_builder *b, nir_ssa_def *instance_addr, nir_ssa_de
 {
    unsigned offset = offsetof(struct radv_bvh_instance_node, wto_matrix);
    for (unsigned i = 0; i < 3; ++i) {
-      out[i] = nir_build_load_global(b, 4, 32,
-                                     nir_iadd(b, instance_addr, nir_imm_int64(b, offset + i * 16)),
+      out[i] = nir_build_load_global(b, 4, 32, nir_iadd_imm(b, instance_addr, offset + i * 16),
                                      .align_mul = 64, .align_offset = offset + i * 16);
    }
 }
@@ -400,28 +399,22 @@ nir_ssa_def *
 hit_is_opaque(nir_builder *b, nir_ssa_def *sbt_offset_and_flags, nir_ssa_def *flags,
               nir_ssa_def *geometry_id_and_flags)
 {
-   nir_ssa_def *geom_force_opaque = nir_ine(
-      b, nir_iand(b, geometry_id_and_flags, nir_imm_int(b, VK_GEOMETRY_OPAQUE_BIT_KHR << 28)),
-      nir_imm_int(b, 0));
-   nir_ssa_def *instance_force_opaque =
-      nir_ine(b,
-              nir_iand(b, sbt_offset_and_flags,
-                       nir_imm_int(b, VK_GEOMETRY_INSTANCE_FORCE_OPAQUE_BIT_KHR << 24)),
-              nir_imm_int(b, 0));
-   nir_ssa_def *instance_force_non_opaque =
-      nir_ine(b,
-              nir_iand(b, sbt_offset_and_flags,
-                       nir_imm_int(b, VK_GEOMETRY_INSTANCE_FORCE_NO_OPAQUE_BIT_KHR << 24)),
-              nir_imm_int(b, 0));
+   nir_ssa_def *geom_force_opaque =
+      nir_ine_imm(b, nir_iand_imm(b, geometry_id_and_flags, VK_GEOMETRY_OPAQUE_BIT_KHR << 28), 0);
+   nir_ssa_def *instance_force_opaque = nir_ine_imm(
+      b, nir_iand_imm(b, sbt_offset_and_flags, VK_GEOMETRY_INSTANCE_FORCE_OPAQUE_BIT_KHR << 24), 0);
+   nir_ssa_def *instance_force_non_opaque = nir_ine_imm(
+      b, nir_iand_imm(b, sbt_offset_and_flags, VK_GEOMETRY_INSTANCE_FORCE_NO_OPAQUE_BIT_KHR << 24),
+      0);
 
    nir_ssa_def *opaque = geom_force_opaque;
    opaque = nir_bcsel(b, instance_force_opaque, nir_imm_bool(b, true), opaque);
    opaque = nir_bcsel(b, instance_force_non_opaque, nir_imm_bool(b, false), opaque);
 
    nir_ssa_def *ray_force_opaque =
-      nir_ine(b, nir_iand(b, flags, nir_imm_int(b, SpvRayFlagsOpaqueKHRMask)), nir_imm_int(b, 0));
+      nir_ine_imm(b, nir_iand_imm(b, flags, SpvRayFlagsOpaqueKHRMask), 0);
    nir_ssa_def *ray_force_non_opaque =
-      nir_ine(b, nir_iand(b, flags, nir_imm_int(b, SpvRayFlagsNoOpaqueKHRMask)), nir_imm_int(b, 0));
+      nir_ine_imm(b, nir_iand_imm(b, flags, SpvRayFlagsNoOpaqueKHRMask), 0);
 
    opaque = nir_bcsel(b, ray_force_opaque, nir_imm_bool(b, true), opaque);
    opaque = nir_bcsel(b, ray_force_non_opaque, nir_imm_bool(b, false), opaque);
diff --git a/src/amd/vulkan/radv_shader.c b/src/amd/vulkan/radv_shader.c
index eb29657e603..0be40340d7b 100644
--- a/src/amd/vulkan/radv_shader.c
+++ b/src/amd/vulkan/radv_shader.c
@@ -365,12 +365,12 @@ radv_lower_primitive_shading_rate(nir_shader *nir)
          nir_ssa_def *val = nir_ssa_for_src(&b, intr->src[1], 1);
 
          /* x_rate = (shadingRate & (Horizontal2Pixels | Horizontal4Pixels)) ? 0x1 : 0x0; */
-         nir_ssa_def *x_rate = nir_iand(&b, val, nir_imm_int(&b, 12));
-         x_rate = nir_b2i32(&b, nir_ine(&b, x_rate, nir_imm_int(&b, 0)));
+         nir_ssa_def *x_rate = nir_iand_imm(&b, val, 12);
+         x_rate = nir_b2i32(&b, nir_ine_imm(&b, x_rate, 0));
 
          /* y_rate = (shadingRate & (Vertical2Pixels | Vertical4Pixels)) ? 0x1 : 0x0; */
-         nir_ssa_def *y_rate = nir_iand(&b, val, nir_imm_int(&b, 3));
-         y_rate = nir_b2i32(&b, nir_ine(&b, y_rate, nir_imm_int(&b, 0)));
+         nir_ssa_def *y_rate = nir_iand_imm(&b, val, 3);
+         y_rate = nir_b2i32(&b, nir_ine_imm(&b, y_rate, 0));
 
          nir_ssa_def *out = NULL;
 
@@ -383,8 +383,7 @@ radv_lower_primitive_shading_rate(nir_shader *nir)
              * Bits [30:31] = VRS rate Y
              * This will be added to the other bits of that channel in the backend.
              */
-            out = nir_ior(&b, nir_ishl(&b, x_rate, nir_imm_int(&b, 28)),
-                              nir_ishl(&b, y_rate, nir_imm_int(&b, 30)));
+            out = nir_ior(&b, nir_ishl_imm(&b, x_rate, 28), nir_ishl_imm(&b, y_rate, 30));
          } else {
             /* VS, TES, GS:
              * Primitive shading rate is a per-vertex output pos export.
@@ -393,8 +392,7 @@ radv_lower_primitive_shading_rate(nir_shader *nir)
              * Bits [4:5] = VRS rate Y
              * HW shading rate = (xRate << 2) | (yRate << 4)
              */
-            out = nir_ior(&b, nir_ishl(&b, x_rate, nir_imm_int(&b, 2)),
-                              nir_ishl(&b, y_rate, nir_imm_int(&b, 4)));
+            out = nir_ior(&b, nir_ishl_imm(&b, x_rate, 2), nir_ishl_imm(&b, y_rate, 4));
          }
 
          nir_instr_rewrite_src(&intr->instr, &intr->src[1], nir_src_for_ssa(out));
@@ -531,10 +529,10 @@ radv_lower_fs_intrinsics(nir_shader *nir, const struct radv_pipeline_stage *fs_s
             /* VRS Rate X = Ancillary[2:3] */
             nir_ssa_def *ancillary =
                nir_load_vector_arg_amd(&b, 1, .base = args->ac.ancillary.arg_index);
-            nir_ssa_def *x_rate = nir_ubfe(&b, ancillary, nir_imm_int(&b, 2), nir_imm_int(&b, 2));
+            nir_ssa_def *x_rate = nir_ubfe_imm(&b, ancillary, 2, 2);
 
             /* xRate = xRate == 0x1 ? adjusted_frag_z : frag_z. */
-            nir_ssa_def *cond = nir_ieq(&b, x_rate, nir_imm_int(&b, 1));
+            nir_ssa_def *cond = nir_ieq_imm(&b, x_rate, 1);
             frag_z = nir_bcsel(&b, cond, adjusted_frag_z, frag_z);
 
             nir_ssa_def *new_dest = nir_vector_insert_imm(&b, &intrin->dest.ssa, frag_z, 2);