ac/nir: remove nir_to_llvm_context::builder

Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>

diff --git a/src/amd/common/ac_nir_to_llvm.c b/src/amd/common/ac_nir_to_llvm.c
index 09f47fe8d0909126c9cd6a00b8c03fd3e639942c..a052a7109d4ce5faf0ed2c8d7e98d7a28fdbf956 100644 (file)
--- a/src/amd/common/ac_nir_to_llvm.c
+++ b/src/amd/common/ac_nir_to_llvm.c
@@ -80,7 +80,6 @@ struct nir_to_llvm_context {
        unsigned max_workgroup_size;
        LLVMContextRef context;
        LLVMModuleRef module;
-       LLVMBuilderRef builder;
        LLVMValueRef main_function;
 
        LLVMValueRef descriptor_sets[AC_UD_MAX_SETS];
@@ -395,7 +394,7 @@ get_tcs_out_patch_stride(struct nir_to_llvm_context *ctx)
 static LLVMValueRef
 get_tcs_out_patch0_offset(struct nir_to_llvm_context *ctx)
 {
-       return LLVMBuildMul(ctx->builder,
+       return LLVMBuildMul(ctx->ac.builder,
                            unpack_param(&ctx->ac, ctx->tcs_out_offsets, 0, 16),
                            LLVMConstInt(ctx->ac.i32, 4, false), "");
 }
@@ -403,7 +402,7 @@ get_tcs_out_patch0_offset(struct nir_to_llvm_context *ctx)
 static LLVMValueRef
 get_tcs_out_patch0_patch_data_offset(struct nir_to_llvm_context *ctx)
 {
-       return LLVMBuildMul(ctx->builder,
+       return LLVMBuildMul(ctx->ac.builder,
                            unpack_param(&ctx->ac, ctx->tcs_out_offsets, 16, 16),
                            LLVMConstInt(ctx->ac.i32, 4, false), "");
 }
@@ -414,7 +413,7 @@ get_tcs_in_current_patch_offset(struct nir_to_llvm_context *ctx)
        LLVMValueRef patch_stride = get_tcs_in_patch_stride(ctx);
        LLVMValueRef rel_patch_id = get_rel_patch_id(ctx);
 
-       return LLVMBuildMul(ctx->builder, patch_stride, rel_patch_id, "");
+       return LLVMBuildMul(ctx->ac.builder, patch_stride, rel_patch_id, "");
 }
 
 static LLVMValueRef
@@ -424,8 +423,8 @@ get_tcs_out_current_patch_offset(struct nir_to_llvm_context *ctx)
        LLVMValueRef patch_stride = get_tcs_out_patch_stride(ctx);
        LLVMValueRef rel_patch_id = get_rel_patch_id(ctx);
 
-       return LLVMBuildAdd(ctx->builder, patch0_offset,
-                           LLVMBuildMul(ctx->builder, patch_stride,
+       return LLVMBuildAdd(ctx->ac.builder, patch0_offset,
+                           LLVMBuildMul(ctx->ac.builder, patch_stride,
                                         rel_patch_id, ""),
                            "");
 }
@@ -438,8 +437,8 @@ get_tcs_out_current_patch_data_offset(struct nir_to_llvm_context *ctx)
        LLVMValueRef patch_stride = get_tcs_out_patch_stride(ctx);
        LLVMValueRef rel_patch_id = get_rel_patch_id(ctx);
 
-       return LLVMBuildAdd(ctx->builder, patch0_patch_data_offset,
-                           LLVMBuildMul(ctx->builder, patch_stride,
+       return LLVMBuildAdd(ctx->ac.builder, patch0_patch_data_offset,
+                           LLVMBuildMul(ctx->ac.builder, patch_stride,
                                         rel_patch_id, ""),
                            "");
 }
@@ -1021,7 +1020,7 @@ static void create_function(struct nir_to_llvm_context *ctx,
        }
 
        ctx->main_function = create_llvm_function(
-           ctx->context, ctx->module, ctx->builder, NULL, 0, &args,
+           ctx->context, ctx->module, ctx->ac.builder, NULL, 0, &args,
            ctx->max_workgroup_size,
            ctx->options->unsafe_math);
        set_llvm_calling_convention(ctx->main_function, stage);
@@ -1046,7 +1045,7 @@ static void create_function(struct nir_to_llvm_context *ctx,
                        ctx->ring_offsets = ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.implicit.buffer.ptr",
                                                               LLVMPointerType(ctx->ac.i8, AC_CONST_ADDR_SPACE),
                                                               NULL, 0, AC_FUNC_ATTR_READNONE);
-                       ctx->ring_offsets = LLVMBuildBitCast(ctx->builder, ctx->ring_offsets,
+                       ctx->ring_offsets = LLVMBuildBitCast(ctx->ac.builder, ctx->ring_offsets,
                                                             ac_array_in_const_addr_space(ctx->ac.v4i32), "");
                }
        }
@@ -2372,8 +2371,8 @@ radv_load_resource(struct ac_shader_abi *abi, LLVMValueRef index,
                stride = LLVMConstInt(ctx->ac.i32, layout->binding[binding].size, false);
 
        offset = LLVMConstInt(ctx->ac.i32, base_offset, false);
-       index = LLVMBuildMul(ctx->builder, index, stride, "");
-       offset = LLVMBuildAdd(ctx->builder, offset, index, "");
+       index = LLVMBuildMul(ctx->ac.builder, index, stride, "");
+       offset = LLVMBuildAdd(ctx->ac.builder, offset, index, "");
        
        desc_ptr = ac_build_gep0(&ctx->ac, desc_ptr, offset);
        desc_ptr = cast_ptr(&ctx->ac, desc_ptr, ctx->ac.v4i32);
@@ -2771,15 +2770,15 @@ static LLVMValueRef get_tcs_tes_buffer_address(struct nir_to_llvm_context *ctx,
 
        vertices_per_patch = unpack_param(&ctx->ac, ctx->tcs_offchip_layout, 9, 6);
        num_patches = unpack_param(&ctx->ac, ctx->tcs_offchip_layout, 0, 9);
-       total_vertices = LLVMBuildMul(ctx->builder, vertices_per_patch,
+       total_vertices = LLVMBuildMul(ctx->ac.builder, vertices_per_patch,
                                      num_patches, "");
 
        constant16 = LLVMConstInt(ctx->ac.i32, 16, false);
        if (vertex_index) {
-               base_addr = LLVMBuildMul(ctx->builder, rel_patch_id,
+               base_addr = LLVMBuildMul(ctx->ac.builder, rel_patch_id,
                                         vertices_per_patch, "");
 
-               base_addr = LLVMBuildAdd(ctx->builder, base_addr,
+               base_addr = LLVMBuildAdd(ctx->ac.builder, base_addr,
                                         vertex_index, "");
 
                param_stride = total_vertices;
@@ -2788,17 +2787,17 @@ static LLVMValueRef get_tcs_tes_buffer_address(struct nir_to_llvm_context *ctx,
                param_stride = num_patches;
        }
 
-       base_addr = LLVMBuildAdd(ctx->builder, base_addr,
-                                LLVMBuildMul(ctx->builder, param_index,
+       base_addr = LLVMBuildAdd(ctx->ac.builder, base_addr,
+                                LLVMBuildMul(ctx->ac.builder, param_index,
                                              param_stride, ""), "");
 
-       base_addr = LLVMBuildMul(ctx->builder, base_addr, constant16, "");
+       base_addr = LLVMBuildMul(ctx->ac.builder, base_addr, constant16, "");
 
        if (!vertex_index) {
                LLVMValueRef patch_data_offset =
                           unpack_param(&ctx->ac, ctx->tcs_offchip_layout, 16, 16);
 
-               base_addr = LLVMBuildAdd(ctx->builder, base_addr,
+               base_addr = LLVMBuildAdd(ctx->ac.builder, base_addr,
                                         patch_data_offset, "");
        }
        return base_addr;
@@ -2814,7 +2813,7 @@ static LLVMValueRef get_tcs_tes_buffer_address_params(struct nir_to_llvm_context
        LLVMValueRef param_index;
 
        if (indir_index)
-               param_index = LLVMBuildAdd(ctx->builder, LLVMConstInt(ctx->ac.i32, param, false),
+               param_index = LLVMBuildAdd(ctx->ac.builder, LLVMConstInt(ctx->ac.i32, param, false),
                                           indir_index, "");
        else {
                if (const_index && !is_compact)
@@ -2848,25 +2847,25 @@ get_dw_address(struct nir_to_llvm_context *ctx,
 {
 
        if (vertex_index) {
-               dw_addr = LLVMBuildAdd(ctx->builder, dw_addr,
-                                      LLVMBuildMul(ctx->builder,
+               dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
+                                      LLVMBuildMul(ctx->ac.builder,
                                                    vertex_index,
                                                    stride, ""), "");
        }
 
        if (indir_index)
-               dw_addr = LLVMBuildAdd(ctx->builder, dw_addr,
-                                      LLVMBuildMul(ctx->builder, indir_index,
+               dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
+                                      LLVMBuildMul(ctx->ac.builder, indir_index,
                                                    LLVMConstInt(ctx->ac.i32, 4, false), ""), "");
        else if (const_index && !compact_const_index)
-               dw_addr = LLVMBuildAdd(ctx->builder, dw_addr,
+               dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
                                       LLVMConstInt(ctx->ac.i32, const_index, false), "");
 
-       dw_addr = LLVMBuildAdd(ctx->builder, dw_addr,
+       dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
                               LLVMConstInt(ctx->ac.i32, param * 4, false), "");
 
        if (const_index && compact_const_index)
-               dw_addr = LLVMBuildAdd(ctx->builder, dw_addr,
+               dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
                                       LLVMConstInt(ctx->ac.i32, const_index, false), "");
        return dw_addr;
 }
@@ -2907,7 +2906,7 @@ load_tcs_varyings(struct ac_shader_abi *abi,
 
        for (unsigned i = 0; i < num_components + component; i++) {
                value[i] = ac_lds_load(&ctx->ac, dw_addr);
-               dw_addr = LLVMBuildAdd(ctx->builder, dw_addr,
+               dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
                                       ctx->ac.i32_1, "");
        }
        result = ac_build_varying_gather_values(&ctx->ac, value, num_components, component);
@@ -2976,7 +2975,7 @@ store_tcs_output(struct ac_shader_abi *abi,
 
                if (store_lds || is_tess_factor) {
                        LLVMValueRef dw_addr_chan =
-                               LLVMBuildAdd(ctx->builder, dw_addr,
+                               LLVMBuildAdd(ctx->ac.builder, dw_addr,
                                                           LLVMConstInt(ctx->ac.i32, chan, false), "");
                        ac_lds_store(&ctx->ac, dw_addr_chan, value);
                }
@@ -3021,7 +3020,7 @@ load_tes_input(struct ac_shader_abi *abi,
                                                     is_compact, vertex_index, param_index);
 
        LLVMValueRef comp_offset = LLVMConstInt(ctx->ac.i32, component * 4, false);
-       buf_addr = LLVMBuildAdd(ctx->builder, buf_addr, comp_offset, "");
+       buf_addr = LLVMBuildAdd(ctx->ac.builder, buf_addr, comp_offset, "");
 
        result = ac_build_buffer_load(&ctx->ac, ctx->hs_ring_tess_offchip, num_components, NULL,
                                      buf_addr, ctx->oc_lds, is_compact ? (4 * const_index) : 0, 1, 0, true, false);
@@ -3046,7 +3045,7 @@ load_gs_input(struct ac_shader_abi *abi,
 
        vtx_offset_param = vertex_index;
        assert(vtx_offset_param < 6);
-       vtx_offset = LLVMBuildMul(ctx->builder, ctx->gs_vtx_offset[vtx_offset_param],
+       vtx_offset = LLVMBuildMul(ctx->ac.builder, ctx->gs_vtx_offset[vtx_offset_param],
                                  LLVMConstInt(ctx->ac.i32, 4, false), "");
 
        param = shader_io_get_unique_index(location);
@@ -3069,7 +3068,7 @@ load_gs_input(struct ac_shader_abi *abi,
                                                        vtx_offset, soffset,
                                                        0, 1, 0, true, false);
 
-                       value[i] = LLVMBuildBitCast(ctx->builder, value[i],
+                       value[i] = LLVMBuildBitCast(ctx->ac.builder, value[i],
                                                    type, "");
                }
        }
@@ -4003,10 +4002,10 @@ static LLVMValueRef load_sample_position(struct ac_shader_abi *abi,
        LLVMValueRef result;
        LLVMValueRef ptr = ac_build_gep0(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->ac.i32, RING_PS_SAMPLE_POSITIONS, false));
 
-       ptr = LLVMBuildBitCast(ctx->builder, ptr,
+       ptr = LLVMBuildBitCast(ctx->ac.builder, ptr,
                               ac_array_in_const_addr_space(ctx->ac.v2f32), "");
 
-       sample_id = LLVMBuildAdd(ctx->builder, sample_id, ctx->sample_pos_offset, "");
+       sample_id = LLVMBuildAdd(ctx->ac.builder, sample_id, ctx->sample_pos_offset, "");
        result = ac_build_load_invariant(&ctx->ac, ptr, sample_id);
 
        return result;
@@ -4165,7 +4164,7 @@ visit_emit_vertex(struct ac_shader_abi *abi, unsigned stream, LLVMValueRef *addr
        assert(stream == 0);
 
        /* Write vertex attribute values to GSVS ring */
-       gs_next_vertex = LLVMBuildLoad(ctx->builder,
+       gs_next_vertex = LLVMBuildLoad(ctx->ac.builder,
                                       ctx->gs_next_vertex,
                                       "");
 
@@ -4174,7 +4173,7 @@ visit_emit_vertex(struct ac_shader_abi *abi, unsigned stream, LLVMValueRef *addr
         * have any effect, and GS threads have no externally observable
         * effects other than emitting vertices.
         */
-       can_emit = LLVMBuildICmp(ctx->builder, LLVMIntULT, gs_next_vertex,
+       can_emit = LLVMBuildICmp(ctx->ac.builder, LLVMIntULT, gs_next_vertex,
                                 LLVMConstInt(ctx->ac.i32, ctx->gs_max_out_vertices, false), "");
        ac_build_kill_if_false(&ctx->ac, can_emit);
 
@@ -4196,13 +4195,13 @@ visit_emit_vertex(struct ac_shader_abi *abi, unsigned stream, LLVMValueRef *addr
                                slot_inc = 2;
                }
                for (unsigned j = 0; j < length; j++) {
-                       LLVMValueRef out_val = LLVMBuildLoad(ctx->builder,
+                       LLVMValueRef out_val = LLVMBuildLoad(ctx->ac.builder,
                                                             out_ptr[j], "");
                        LLVMValueRef voffset = LLVMConstInt(ctx->ac.i32, (slot * 4 + j) * ctx->gs_max_out_vertices, false);
-                       voffset = LLVMBuildAdd(ctx->builder, voffset, gs_next_vertex, "");
-                       voffset = LLVMBuildMul(ctx->builder, voffset, LLVMConstInt(ctx->ac.i32, 4, false), "");
+                       voffset = LLVMBuildAdd(ctx->ac.builder, voffset, gs_next_vertex, "");
+                       voffset = LLVMBuildMul(ctx->ac.builder, voffset, LLVMConstInt(ctx->ac.i32, 4, false), "");
 
-                       out_val = LLVMBuildBitCast(ctx->builder, out_val, ctx->ac.i32, "");
+                       out_val = LLVMBuildBitCast(ctx->ac.builder, out_val, ctx->ac.i32, "");
 
                        ac_build_buffer_store_dword(&ctx->ac, ctx->gsvs_ring,
                                                    out_val, 1,
@@ -4212,9 +4211,9 @@ visit_emit_vertex(struct ac_shader_abi *abi, unsigned stream, LLVMValueRef *addr
                idx += slot_inc;
        }
 
-       gs_next_vertex = LLVMBuildAdd(ctx->builder, gs_next_vertex,
+       gs_next_vertex = LLVMBuildAdd(ctx->ac.builder, gs_next_vertex,
                                      ctx->ac.i32_1, "");
-       LLVMBuildStore(ctx->builder, gs_next_vertex, ctx->gs_next_vertex);
+       LLVMBuildStore(ctx->ac.builder, gs_next_vertex, ctx->gs_next_vertex);
 
        ac_build_sendmsg(&ctx->ac, AC_SENDMSG_GS_OP_EMIT | AC_SENDMSG_GS | (0 << 8), ctx->gs_wave_id);
 }
@@ -4239,8 +4238,8 @@ load_tess_coord(struct ac_shader_abi *abi)
        };
 
        if (ctx->tes_primitive_mode == GL_TRIANGLES)
-               coord[2] = LLVMBuildFSub(ctx->builder, ctx->ac.f32_1,
-                                       LLVMBuildFAdd(ctx->builder, coord[0], coord[1], ""), "");
+               coord[2] = LLVMBuildFSub(ctx->ac.builder, ctx->ac.f32_1,
+                                       LLVMBuildFAdd(ctx->ac.builder, coord[0], coord[1], ""), "");
 
        return ac_build_gather_values(&ctx->ac, coord, 3);
 }
@@ -4531,7 +4530,7 @@ static LLVMValueRef radv_load_ssbo(struct ac_shader_abi *abi,
 
        LLVMSetMetadata(buffer_ptr, ctx->ac.uniform_md_kind, ctx->ac.empty_md);
 
-       result = LLVMBuildLoad(ctx->builder, buffer_ptr, "");
+       result = LLVMBuildLoad(ctx->ac.builder, buffer_ptr, "");
        LLVMSetMetadata(result, ctx->ac.invariant_load_md_kind, ctx->ac.empty_md);
 
        return result;
@@ -4544,7 +4543,7 @@ static LLVMValueRef radv_load_ubo(struct ac_shader_abi *abi, LLVMValueRef buffer
 
        LLVMSetMetadata(buffer_ptr, ctx->ac.uniform_md_kind, ctx->ac.empty_md);
 
-       result = LLVMBuildLoad(ctx->builder, buffer_ptr, "");
+       result = LLVMBuildLoad(ctx->ac.builder, buffer_ptr, "");
        LLVMSetMetadata(result, ctx->ac.invariant_load_md_kind, ctx->ac.empty_md);
 
        return result;
@@ -4565,7 +4564,7 @@ static LLVMValueRef radv_get_sampler_desc(struct ac_shader_abi *abi,
        unsigned offset = binding->offset;
        unsigned stride = binding->size;
        unsigned type_size;
-       LLVMBuilderRef builder = ctx->builder;
+       LLVMBuilderRef builder = ctx->ac.builder;
        LLVMTypeRef type;
 
        assert(base_index < layout->binding_count);
@@ -5327,7 +5326,7 @@ handle_vs_input_decl(struct nir_to_llvm_context *ctx,
 
        for (unsigned i = 0; i < attrib_count; ++i, ++idx) {
                if (ctx->options->key.vs.instance_rate_inputs & (1u << (index + i))) {
-                       buffer_index = LLVMBuildAdd(ctx->builder, ctx->abi.instance_id,
+                       buffer_index = LLVMBuildAdd(ctx->ac.builder, ctx->abi.instance_id,
                                                    ctx->abi.start_instance, "");
                        if (ctx->options->key.vs.as_ls) {
                                ctx->shader_info->vs.vgpr_comp_cnt =
@@ -5337,7 +5336,7 @@ handle_vs_input_decl(struct nir_to_llvm_context *ctx,
                                        MAX2(1, ctx->shader_info->vs.vgpr_comp_cnt);
                        }
                } else
-                       buffer_index = LLVMBuildAdd(ctx->builder, ctx->abi.vertex_id,
+                       buffer_index = LLVMBuildAdd(ctx->ac.builder, ctx->abi.vertex_id,
                                                    ctx->abi.base_vertex, "");
                t_offset = LLVMConstInt(ctx->ac.i32, index + i, false);
 
@@ -5353,7 +5352,7 @@ handle_vs_input_decl(struct nir_to_llvm_context *ctx,
                for (unsigned chan = 0; chan < 4; chan++) {
                        LLVMValueRef llvm_chan = LLVMConstInt(ctx->ac.i32, chan, false);
                        ctx->inputs[radeon_llvm_reg_index_soa(idx, chan)] =
-                               ac_to_integer(&ctx->ac, LLVMBuildExtractElement(ctx->builder,
+                               ac_to_integer(&ctx->ac, LLVMBuildExtractElement(ctx->ac.builder,
                                                        input, llvm_chan, ""));
                }
        }
@@ -5383,12 +5382,12 @@ static void interp_fs_input(struct nir_to_llvm_context *ctx,
         * to NaN.
         */
        if (interp) {
-               interp_param = LLVMBuildBitCast(ctx->builder, interp_param,
+               interp_param = LLVMBuildBitCast(ctx->ac.builder, interp_param,
                                                ctx->ac.v2f32, "");
 
-               i = LLVMBuildExtractElement(ctx->builder, interp_param,
+               i = LLVMBuildExtractElement(ctx->ac.builder, interp_param,
                                                ctx->ac.i32_0, "");
-               j = LLVMBuildExtractElement(ctx->builder, interp_param,
+               j = LLVMBuildExtractElement(ctx->ac.builder, interp_param,
                                                ctx->ac.i32_1, "");
        }
 
@@ -5468,9 +5467,9 @@ prepare_interp_optimize(struct nir_to_llvm_context *ctx,
        }
 
        if (uses_center && uses_centroid) {
-               LLVMValueRef sel = LLVMBuildICmp(ctx->builder, LLVMIntSLT, ctx->abi.prim_mask, ctx->ac.i32_0, "");
-               ctx->persp_centroid = LLVMBuildSelect(ctx->builder, sel, ctx->persp_center, ctx->persp_centroid, "");
-               ctx->linear_centroid = LLVMBuildSelect(ctx->builder, sel, ctx->linear_center, ctx->linear_centroid, "");
+               LLVMValueRef sel = LLVMBuildICmp(ctx->ac.builder, LLVMIntSLT, ctx->abi.prim_mask, ctx->ac.i32_0, "");
+               ctx->persp_centroid = LLVMBuildSelect(ctx->ac.builder, sel, ctx->persp_center, ctx->persp_centroid, "");
+               ctx->linear_centroid = LLVMBuildSelect(ctx->ac.builder, sel, ctx->linear_center, ctx->linear_centroid, "");
        }
 }
 
@@ -5882,7 +5881,7 @@ radv_load_output(struct nir_to_llvm_context *ctx, unsigned index, unsigned chan)
        LLVMValueRef output =
                ctx->nir->outputs[radeon_llvm_reg_index_soa(index, chan)];
 
-       return LLVMBuildLoad(ctx->builder, output, "");
+       return LLVMBuildLoad(ctx->ac.builder, output, "");
 }
 
 static void
@@ -5905,7 +5904,7 @@ handle_vs_outputs_post(struct nir_to_llvm_context *ctx,
                                            si_build_alloca_undef(&ctx->ac, ctx->ac.f32, "");
                }
 
-               LLVMBuildStore(ctx->builder, ac_to_float(&ctx->ac, ctx->abi.view_index),  *tmp_out);
+               LLVMBuildStore(ctx->ac.builder, ac_to_float(&ctx->ac, ctx->abi.view_index),  *tmp_out);
                ctx->output_mask |= 1ull << VARYING_SLOT_LAYER;
        }
 
@@ -5987,10 +5986,10 @@ handle_vs_outputs_post(struct nir_to_llvm_context *ctx,
                                 */
                                LLVMValueRef v = viewport_index_value;
                                v = ac_to_integer(&ctx->ac, v);
-                               v = LLVMBuildShl(ctx->builder, v,
+                               v = LLVMBuildShl(ctx->ac.builder, v,
                                                 LLVMConstInt(ctx->ac.i32, 16, false),
                                                 "");
-                               v = LLVMBuildOr(ctx->builder, v,
+                               v = LLVMBuildOr(ctx->ac.builder, v,
                                                ac_to_integer(&ctx->ac, pos_args[1].out[2]), "");
 
                                pos_args[1].out[2] = ac_to_float(&ctx->ac, v);
@@ -6108,18 +6107,18 @@ handle_es_outputs_post(struct nir_to_llvm_context *ctx,
                param_index = shader_io_get_unique_index(i);
 
                if (lds_base) {
-                       dw_addr = LLVMBuildAdd(ctx->builder, lds_base,
+                       dw_addr = LLVMBuildAdd(ctx->ac.builder, lds_base,
                                               LLVMConstInt(ctx->ac.i32, param_index * 4, false),
                                               "");
                }
                for (j = 0; j < length; j++) {
-                       LLVMValueRef out_val = LLVMBuildLoad(ctx->builder, out_ptr[j], "");
-                       out_val = LLVMBuildBitCast(ctx->builder, out_val, ctx->ac.i32, "");
+                       LLVMValueRef out_val = LLVMBuildLoad(ctx->ac.builder, out_ptr[j], "");
+                       out_val = LLVMBuildBitCast(ctx->ac.builder, out_val, ctx->ac.i32, "");
 
                        if (ctx->ac.chip_class  >= GFX9) {
                                ac_lds_store(&ctx->ac, dw_addr,
-                                            LLVMBuildLoad(ctx->builder, out_ptr[j], ""));
-                               dw_addr = LLVMBuildAdd(ctx->builder, dw_addr, ctx->ac.i32_1, "");
+                                            LLVMBuildLoad(ctx->ac.builder, out_ptr[j], ""));
+                               dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr, ctx->ac.i32_1, "");
                        } else {
                                ac_build_buffer_store_dword(&ctx->ac,
                                                            ctx->esgs_ring,
@@ -6137,7 +6136,7 @@ handle_ls_outputs_post(struct nir_to_llvm_context *ctx)
 {
        LLVMValueRef vertex_id = ctx->rel_auto_id;
        LLVMValueRef vertex_dw_stride = unpack_param(&ctx->ac, ctx->ls_out_layout, 13, 8);
-       LLVMValueRef base_dw_addr = LLVMBuildMul(ctx->builder, vertex_id,
+       LLVMValueRef base_dw_addr = LLVMBuildMul(ctx->ac.builder, vertex_id,
                                                 vertex_dw_stride, "");
 
        for (unsigned i = 0; i < RADEON_LLVM_MAX_OUTPUTS; ++i) {
@@ -6153,13 +6152,13 @@ handle_ls_outputs_post(struct nir_to_llvm_context *ctx)
                mark_tess_output(ctx, false, param);
                if (length > 4)
                        mark_tess_output(ctx, false, param + 1);
-               LLVMValueRef dw_addr = LLVMBuildAdd(ctx->builder, base_dw_addr,
+               LLVMValueRef dw_addr = LLVMBuildAdd(ctx->ac.builder, base_dw_addr,
                                                    LLVMConstInt(ctx->ac.i32, param * 4, false),
                                                    "");
                for (unsigned j = 0; j < length; j++) {
                        ac_lds_store(&ctx->ac, dw_addr,
-                                    LLVMBuildLoad(ctx->builder, out_ptr[j], ""));
-                       dw_addr = LLVMBuildAdd(ctx->builder, dw_addr, ctx->ac.i32_1, "");
+                                    LLVMBuildLoad(ctx->ac.builder, out_ptr[j], ""));
+                       dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr, ctx->ac.i32_1, "");
                }
        }
 }
@@ -6182,7 +6181,7 @@ ac_build_insert_new_block(struct nir_to_llvm_context *ctx, const char *name)
        LLVMBasicBlockRef new_block;
 
        /* get current basic block */
-       current_block = LLVMGetInsertBlock(ctx->builder);
+       current_block = LLVMGetInsertBlock(ctx->ac.builder);
 
        /* chqeck if there's another block after this one */
        next_block = LLVMGetNextBasicBlock(current_block);
@@ -6203,7 +6202,7 @@ ac_nir_build_if(struct ac_build_if_state *ifthen,
                struct nir_to_llvm_context *ctx,
                LLVMValueRef condition)
 {
-       LLVMBasicBlockRef block = LLVMGetInsertBlock(ctx->builder);
+       LLVMBasicBlockRef block = LLVMGetInsertBlock(ctx->ac.builder);
 
        memset(ifthen, 0, sizeof *ifthen);
        ifthen->ctx = ctx;
@@ -6220,7 +6219,7 @@ ac_nir_build_if(struct ac_build_if_state *ifthen,
                                              "if-true-block");
 
        /* successive code goes into the true block */
-       LLVMPositionBuilderAtEnd(ctx->builder, ifthen->true_block);
+       LLVMPositionBuilderAtEnd(ctx->ac.builder, ifthen->true_block);
 }
 
 /**
@@ -6229,7 +6228,7 @@ ac_nir_build_if(struct ac_build_if_state *ifthen,
 static void
 ac_nir_build_endif(struct ac_build_if_state *ifthen)
 {
-       LLVMBuilderRef builder = ifthen->ctx->builder;
+       LLVMBuilderRef builder = ifthen->ctx->ac.builder;
 
        /* Insert branch to the merge block from current block */
        LLVMBuildBr(builder, ifthen->merge_block);
@@ -6289,7 +6288,7 @@ write_tess_factors(struct nir_to_llvm_context *ctx)
        }
 
        ac_nir_build_if(&if_ctx, ctx,
-                       LLVMBuildICmp(ctx->builder, LLVMIntEQ,
+                       LLVMBuildICmp(ctx->ac.builder, LLVMIntEQ,
                                      invocation_id, ctx->ac.i32_0, ""));
 
        tess_inner_index = shader_io_get_unique_index(VARYING_SLOT_TESS_LEVEL_INNER);
@@ -6298,9 +6297,9 @@ write_tess_factors(struct nir_to_llvm_context *ctx)
        mark_tess_output(ctx, true, tess_inner_index);
        mark_tess_output(ctx, true, tess_outer_index);
        lds_base = get_tcs_out_current_patch_data_offset(ctx);
-       lds_inner = LLVMBuildAdd(ctx->builder, lds_base,
+       lds_inner = LLVMBuildAdd(ctx->ac.builder, lds_base,
                                 LLVMConstInt(ctx->ac.i32, tess_inner_index * 4, false), "");
-       lds_outer = LLVMBuildAdd(ctx->builder, lds_base,
+       lds_outer = LLVMBuildAdd(ctx->ac.builder, lds_base,
                                 LLVMConstInt(ctx->ac.i32, tess_outer_index * 4, false), "");
 
        for (i = 0; i < 4; i++) {
@@ -6311,20 +6310,20 @@ write_tess_factors(struct nir_to_llvm_context *ctx)
        // LINES reverseal
        if (ctx->options->key.tcs.primitive_mode == GL_ISOLINES) {
                outer[0] = out[1] = ac_lds_load(&ctx->ac, lds_outer);
-               lds_outer = LLVMBuildAdd(ctx->builder, lds_outer,
+               lds_outer = LLVMBuildAdd(ctx->ac.builder, lds_outer,
                                         ctx->ac.i32_1, "");
                outer[1] = out[0] = ac_lds_load(&ctx->ac, lds_outer);
        } else {
                for (i = 0; i < outer_comps; i++) {
                        outer[i] = out[i] =
                                ac_lds_load(&ctx->ac, lds_outer);
-                       lds_outer = LLVMBuildAdd(ctx->builder, lds_outer,
+                       lds_outer = LLVMBuildAdd(ctx->ac.builder, lds_outer,
                                                 ctx->ac.i32_1, "");
                }
                for (i = 0; i < inner_comps; i++) {
                        inner[i] = out[outer_comps+i] =
                                ac_lds_load(&ctx->ac, lds_inner);
-                       lds_inner = LLVMBuildAdd(ctx->builder, lds_inner,
+                       lds_inner = LLVMBuildAdd(ctx->ac.builder, lds_inner,
                                                 ctx->ac.i32_1, "");
                }
        }
@@ -6339,13 +6338,13 @@ write_tess_factors(struct nir_to_llvm_context *ctx)
 
        buffer = ctx->hs_ring_tess_factor;
        tf_base = ctx->tess_factor_offset;
-       byteoffset = LLVMBuildMul(ctx->builder, rel_patch_id,
+       byteoffset = LLVMBuildMul(ctx->ac.builder, rel_patch_id,
                                  LLVMConstInt(ctx->ac.i32, 4 * stride, false), "");
        unsigned tf_offset = 0;
 
        if (ctx->options->chip_class <= VI) {
                ac_nir_build_if(&inner_if_ctx, ctx,
-                               LLVMBuildICmp(ctx->builder, LLVMIntEQ,
+                               LLVMBuildICmp(ctx->ac.builder, LLVMIntEQ,
                                              rel_patch_id, ctx->ac.i32_0, ""));
 
                /* Store the dynamic HS control word. */
@@ -6556,7 +6555,7 @@ static void ac_llvm_finalize_module(struct nir_to_llvm_context * ctx)
        LLVMRunFunctionPassManager(passmgr, ctx->main_function);
        LLVMFinalizeFunctionPassManager(passmgr);
 
-       LLVMDisposeBuilder(ctx->builder);
+       LLVMDisposeBuilder(ctx->ac.builder);
        LLVMDisposePassManager(passmgr);
 }
 
@@ -6609,12 +6608,12 @@ ac_setup_rings(struct nir_to_llvm_context *ctx)
                ctx->esgs_ring = ac_build_load_to_sgpr(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->ac.i32, RING_ESGS_GS, false));
                ctx->gsvs_ring = ac_build_load_to_sgpr(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->ac.i32, RING_GSVS_GS, false));
 
-               ctx->gsvs_ring = LLVMBuildBitCast(ctx->builder, ctx->gsvs_ring, ctx->ac.v4i32, "");
+               ctx->gsvs_ring = LLVMBuildBitCast(ctx->ac.builder, ctx->gsvs_ring, ctx->ac.v4i32, "");
 
-               ctx->gsvs_ring = LLVMBuildInsertElement(ctx->builder, ctx->gsvs_ring, ctx->gsvs_num_entries, LLVMConstInt(ctx->ac.i32, 2, false), "");
-               tmp = LLVMBuildExtractElement(ctx->builder, ctx->gsvs_ring, ctx->ac.i32_1, "");
-               tmp = LLVMBuildOr(ctx->builder, tmp, ctx->gsvs_ring_stride, "");
-               ctx->gsvs_ring = LLVMBuildInsertElement(ctx->builder, ctx->gsvs_ring, tmp, ctx->ac.i32_1, "");
+               ctx->gsvs_ring = LLVMBuildInsertElement(ctx->ac.builder, ctx->gsvs_ring, ctx->gsvs_num_entries, LLVMConstInt(ctx->ac.i32, 2, false), "");
+               tmp = LLVMBuildExtractElement(ctx->ac.builder, ctx->gsvs_ring, ctx->ac.i32_1, "");
+               tmp = LLVMBuildOr(ctx->ac.builder, tmp, ctx->gsvs_ring_stride, "");
+               ctx->gsvs_ring = LLVMBuildInsertElement(ctx->ac.builder, ctx->gsvs_ring, tmp, ctx->ac.i32_1, "");
        }
 
        if (ctx->stage == MESA_SHADER_TESS_CTRL ||
@@ -6751,8 +6750,7 @@ LLVMModuleRef ac_translate_nir_to_llvm(LLVMTargetMachineRef tm,
                options->unsafe_math ? AC_FLOAT_MODE_UNSAFE_FP_MATH :
                                       AC_FLOAT_MODE_DEFAULT;
 
-       ctx.builder = ac_create_builder(ctx.context, float_mode);
-       ctx.ac.builder = ctx.builder;
+       ctx.ac.builder = ac_create_builder(ctx.context, float_mode);
 
        memset(shader_info, 0, sizeof(*shader_info));
 
@@ -6883,7 +6881,7 @@ LLVMModuleRef ac_translate_nir_to_llvm(LLVMTargetMachineRef tm,
                }
        }
 
-       LLVMBuildRetVoid(ctx.builder);
+       LLVMBuildRetVoid(ctx.ac.builder);
 
        if (options->dump_preoptir)
                ac_dump_module(ctx.module);
@@ -7105,7 +7103,7 @@ static void
 ac_gs_copy_shader_emit(struct nir_to_llvm_context *ctx)
 {
        LLVMValueRef vtx_offset =
-               LLVMBuildMul(ctx->builder, ctx->abi.vertex_id,
+               LLVMBuildMul(ctx->ac.builder, ctx->abi.vertex_id,
                             LLVMConstInt(ctx->ac.i32, 4, false), "");
        int idx = 0;
 
@@ -7135,7 +7133,7 @@ ac_gs_copy_shader_emit(struct nir_to_llvm_context *ctx)
                                                     vtx_offset, soffset,
                                                     0, 1, 1, true, false);
 
-                       LLVMBuildStore(ctx->builder,
+                       LLVMBuildStore(ctx->ac.builder,
                                       ac_to_float(&ctx->ac, value), ctx->nir->outputs[radeon_llvm_reg_index_soa(i, j)]);
                }
                idx += slot_inc;
@@ -7168,8 +7166,7 @@ void ac_create_gs_copy_shader(LLVMTargetMachineRef tm,
                options->unsafe_math ? AC_FLOAT_MODE_UNSAFE_FP_MATH :
                                       AC_FLOAT_MODE_DEFAULT;
 
-       ctx.builder = ac_create_builder(ctx.context, float_mode);
-       ctx.ac.builder = ctx.builder;
+       ctx.ac.builder = ac_create_builder(ctx.context, float_mode);
        ctx.stage = MESA_SHADER_VERTEX;
 
        create_function(&ctx, MESA_SHADER_VERTEX, false, MESA_SHADER_VERTEX);
@@ -7196,7 +7193,7 @@ void ac_create_gs_copy_shader(LLVMTargetMachineRef tm,
 
        ctx.nir = NULL;
 
-       LLVMBuildRetVoid(ctx.builder);
+       LLVMBuildRetVoid(ctx.ac.builder);
 
        ac_llvm_finalize_module(&ctx);